diff --git a/projects/rocr-runtime/rocrtst/common/common.cc b/projects/rocr-runtime/rocrtst/common/common.cc old mode 100755 new mode 100644 index 5bc2450e68..89b80f025f --- a/projects/rocr-runtime/rocrtst/common/common.cc +++ b/projects/rocr-runtime/rocrtst/common/common.cc @@ -410,6 +410,10 @@ hsa_status_t AcquirePoolInfo(hsa_amd_memory_pool_t pool, &pool_i->alloc_granule); RET_IF_HSA_COMMON_ERR(err); + err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_REC_GRANULE, + &pool_i->alloc_rec_granule); + RET_IF_HSA_COMMON_ERR(err); + err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALIGNMENT, &pool_i->alloc_alignment); @@ -445,6 +449,9 @@ hsa_status_t DumpMemoryPoolInfo(const pool_info_t *pool_i, fprintf(stdout, "%s%-28s%-36s\n", ind_lvl.c_str(), "Pool Alloc Granule:", gr_str.c_str()); + std::string recgr_str = std::to_string(pool_i->alloc_rec_granule / 1024) + "KB"; + fprintf(stdout, "%s%-28s%-36s\n", ind_lvl.c_str(), + "Pool Alloc Recommended Granule:", recgr_str.c_str()); std::string al_str = std::to_string(pool_i->alloc_alignment / 1024) + "KB"; diff --git a/projects/rocr-runtime/rocrtst/common/common.h b/projects/rocr-runtime/rocrtst/common/common.h old mode 100755 new mode 100644 index 6e8017d3e0..abca31ca75 --- a/projects/rocr-runtime/rocrtst/common/common.h +++ b/projects/rocr-runtime/rocrtst/common/common.h @@ -85,6 +85,7 @@ typedef struct pool_info_t_ { bool alloc_allowed; size_t alloc_granule; size_t alloc_alignment; + size_t alloc_rec_granule; bool accessible_by_all; uint32_t global_flag; uint64_t aggregate_alloc_max; @@ -92,6 +93,7 @@ typedef struct pool_info_t_ { if (a.segment == segment && a.size == size && a.alloc_allowed == alloc_allowed && a.alloc_granule == alloc_granule + && a.alloc_rec_granule == alloc_rec_granule && a.alloc_alignment == alloc_alignment && a.accessible_by_all == accessible_by_all && a.aggregate_alloc_max == aggregate_alloc_max diff --git a/projects/rocr-runtime/rocrtst/common/helper_funcs.h b/projects/rocr-runtime/rocrtst/common/helper_funcs.h index 5e05836668..28e513ccdf 100755 --- a/projects/rocr-runtime/rocrtst/common/helper_funcs.h +++ b/projects/rocr-runtime/rocrtst/common/helper_funcs.h @@ -154,5 +154,6 @@ static __forceinline ScopeGuard MakeScopeGuard(lambda rel) { MAKE_SCOPE_GUARD_HELPER(PASTE(scopeGuardLambda, __COUNTER__), name, \ __VA_ARGS__) +#define ASSERT_SUCCESS(_val) ASSERT_EQ(HSA_STATUS_SUCCESS, (_val)) } // namespace rocrtst #endif // ROCRTST_COMMON_HELPER_FUNCS_H_ diff --git a/projects/rocr-runtime/rocrtst/samples/rocrinfo/rocrinfo.cc b/projects/rocr-runtime/rocrtst/samples/rocrinfo/rocrinfo.cc old mode 100755 new mode 100644 index 321c6769ee..d03048c0bd --- a/projects/rocr-runtime/rocrtst/samples/rocrinfo/rocrinfo.cc +++ b/projects/rocr-runtime/rocrtst/samples/rocrinfo/rocrinfo.cc @@ -105,6 +105,7 @@ struct pool_info_t { size_t pool_size; bool alloc_allowed; size_t alloc_granule; + size_t alloc_recommended_granule; size_t pool_alloc_alignment; bool pl_access; uint32_t global_flag; @@ -509,6 +510,11 @@ static hsa_status_t AcquirePoolInfo(hsa_amd_memory_pool_t pool, &pool_i->pool_alloc_alignment); RET_IF_HSA_ERR(err); + err = + hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_REC_GRANULE, + &pool_i->alloc_recommended_granule); + RET_IF_HSA_ERR(err); + err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_ACCESSIBLE_BY_ALL, &pool_i->pl_access); diff --git a/projects/rocr-runtime/rocrtst/suites/functional/virtual_memory.cc b/projects/rocr-runtime/rocrtst/suites/functional/virtual_memory.cc new file mode 100644 index 0000000000..9cde47a2d1 --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/functional/virtual_memory.cc @@ -0,0 +1,826 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2018, Advanced Micro Devices, Inc. + * All rights reserved. + * + * Developed by: + * + * AMD Research and AMD ROC Software Development + * + * Advanced Micro Devices, Inc. + * + * www.amd.com + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal with the Software without restriction, including without limitation + * the rights to use, copy, modify, merge, publish, distribute, sublicense, + * and/or sell copies of the Software, and to permit persons to whom the + * Software is furnished to do so, subject to the following conditions: + * + * - Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimers. + * - Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimers in + * the documentation and/or other materials provided with the distribution. + * - Neither the names of , + * nor the names of its contributors may be used to endorse or promote + * products derived from this Software without specific prior written + * permission. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL + * THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR + * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, + * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER + * DEALINGS WITH THE SOFTWARE. + * + */ + + +#include +#include +#include +#include +#include +#include +#include + +#include "suites/functional/virtual_memory.h" +#include "common/base_rocr_utils.h" +#include "common/common.h" +#include "common/helper_funcs.h" +#include "common/hsatimer.h" +#include "common/concurrent_utils.h" +#include "gtest/gtest.h" +#include "hsa/hsa.h" + +// Wrap printf to add first or second process indicator +#define PROCESS_LOG(format, ...) \ + { \ + if (verbosity() >= VERBOSE_STANDARD || !parentProcess_) { \ + fprintf(stdout, "line:%d P%u: " format, __LINE__, static_cast(!parentProcess_), \ + ##__VA_ARGS__); \ + } \ + } + +// Fork safe ASSERT_EQ. +#define MSG(y, msg, ...) msg +#define Y(y, ...) y + +#define FORK_ASSERT_EQ(x, ...) \ + if ((x) != (Y(__VA_ARGS__))) { \ + if ((x) != (Y(__VA_ARGS__))) { \ + std::cout << MSG(__VA_ARGS__, ""); \ + if (parentProcess_) { \ + shared_->parent_status = -1; \ + } else { \ + shared_->child_status = -1; \ + } \ + ASSERT_EQ(x, Y(__VA_ARGS__)); \ + } \ + } + +static const char kSubTestSeparator[] = " **************************"; + +static void PrintMemorySubtestHeader(const char* header) { + std::cout << " *** Virtual Memory Functional Subtest: " << header << " ***" << std::endl; +} + +VirtMemoryTestBasic::VirtMemoryTestBasic(void) : TestBase() { + set_title("ROCr Virtual Memory Basic Tests"); + set_description(" Tests virtual memory API functions"); +} + +VirtMemoryTestBasic::~VirtMemoryTestBasic(void) {} + +void VirtMemoryTestBasic::TestCreateDestroy(hsa_agent_t agent, hsa_amd_memory_pool_t pool) { + std::vector gpus; + rocrtst::pool_info_t pool_i; + hsa_device_type_t ag_type; + char ag_name[64]; + void* addrRangeUnmapped; + hsa_status_t err; + void* addrRange; + + ASSERT_SUCCESS(hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &ag_type)); + + ASSERT_SUCCESS(rocrtst::AcquirePoolInfo(pool, &pool_i)); + + if (ag_type != HSA_DEVICE_TYPE_GPU || !pool_i.alloc_allowed) return; + + size_t granule_size = pool_i.alloc_granule; + + ASSERT_SUCCESS(hsa_iterate_agents(rocrtst::IterateGPUAgents, &gpus)); + ASSERT_SUCCESS(hsa_amd_vmem_address_reserve(&addrRange, 20 * granule_size, 0, 0)); + ASSERT_SUCCESS(hsa_amd_vmem_address_reserve(&addrRangeUnmapped, 10 * granule_size, 0, 0)); + + hsa_amd_vmem_alloc_handle_t mem_handle; + ASSERT_SUCCESS( + hsa_amd_vmem_handle_create(pool, 10 * granule_size, MEMORY_TYPE_NONE, 0, &mem_handle)); + + /* Test alloc properties returns correct memory type and pool handle */ + hsa_amd_memory_pool_t poolRet; + hsa_amd_memory_type_t memTypeRet; + ASSERT_SUCCESS(hsa_amd_vmem_get_alloc_properties_from_handle(mem_handle, &poolRet, &memTypeRet)); + + ASSERT_EQ(poolRet.handle, pool.handle); + ASSERT_EQ(memTypeRet, MEMORY_TYPE_NONE); + + hsa_amd_vmem_alloc_handle_t mem_handleTypePinned; + ASSERT_SUCCESS(hsa_amd_vmem_handle_create(pool, 10 * granule_size, MEMORY_TYPE_PINNED, 0, + &mem_handleTypePinned)); + + ASSERT_SUCCESS( + hsa_amd_vmem_get_alloc_properties_from_handle(mem_handleTypePinned, &poolRet, &memTypeRet)); + ASSERT_EQ(poolRet.handle, pool.handle); + ASSERT_EQ(memTypeRet, MEMORY_TYPE_PINNED); + + + ASSERT_SUCCESS(hsa_amd_vmem_map(addrRange, 10 * granule_size, 0, mem_handle, 0)); + + // Access to each GPU should be None + for (auto gpuIt = gpus.begin(); gpuIt != gpus.end(); ++gpuIt) { + hsa_access_permission_t perm = HSA_ACCESS_PERMISSION_RW; + + ASSERT_SUCCESS(hsa_amd_vmem_get_access(addrRange, &perm, *gpuIt)); + ASSERT_EQ(perm, HSA_ACCESS_PERMISSION_NONE); + } + + /* Set RO Access to all GPUs */ + { + int descIndex = 0; + hsa_amd_memory_access_desc_t desc[gpus.size()]; + for (auto gpuIt = gpus.begin(); gpuIt != gpus.end(); ++gpuIt) { + desc[descIndex++] = {HSA_ACCESS_PERMISSION_RO, *gpuIt}; + } + + ASSERT_SUCCESS(hsa_amd_vmem_set_access(addrRange, 10 * granule_size, desc, gpus.size())); + } + + for (auto gpuIt = gpus.begin(); gpuIt != gpus.end(); ++gpuIt) { + hsa_access_permission_t perm = HSA_ACCESS_PERMISSION_NONE; + + ASSERT_SUCCESS(hsa_amd_vmem_get_access(addrRange, &perm, *gpuIt)); + ASSERT_EQ(perm, HSA_ACCESS_PERMISSION_RO); + + /* addrRangeUnmapped was never mapped, so this is an invalid mapping */ + err = hsa_amd_vmem_get_access(addrRangeUnmapped, &perm, *gpuIt); + ASSERT_EQ(err, HSA_STATUS_ERROR_INVALID_ALLOCATION); + } + + if (gpus.size() > 1) { + /* Call set_access with a smaller list of agents, this should remove access for the agents that + * are not included */ + hsa_amd_memory_access_desc_t desc = {HSA_ACCESS_PERMISSION_RW, gpus[1]}; + ASSERT_SUCCESS(hsa_amd_vmem_set_access(addrRange, 10 * granule_size, &desc, 1)); + + size_t i = 0; + for (i = 0; i < gpus.size(); i++) { + hsa_access_permission_t perm = HSA_ACCESS_PERMISSION_NONE; + + /* Only 2nd GPU should have RW access */ + ASSERT_SUCCESS(hsa_amd_vmem_get_access(addrRange, &perm, gpus[i])); + if (i == 1) { + ASSERT_EQ(perm, HSA_ACCESS_PERMISSION_RW); + } else { + ASSERT_EQ(perm, HSA_ACCESS_PERMISSION_NONE); + } + } + } + + ASSERT_SUCCESS(hsa_amd_vmem_unmap(addrRange, 10 * granule_size)); + ASSERT_SUCCESS(hsa_amd_vmem_handle_release(mem_handle)); + ASSERT_SUCCESS(hsa_amd_vmem_address_free(addrRange, 20 * granule_size)); + ASSERT_SUCCESS(hsa_amd_vmem_address_free(addrRangeUnmapped, 10 * granule_size)); +} + +void VirtMemoryTestBasic::TestCreateDestroy(void) { + hsa_status_t err; + std::vector> agent_pools; + + if (verbosity() > 0) { + PrintMemorySubtestHeader("CreateDestroy Test"); + } + bool supp = false; + ASSERT_SUCCESS(hsa_system_get_info(HSA_AMD_SYSTEM_INFO_VIRTUAL_MEM_API_SUPPORTED, (void*)&supp)); + if (!supp) { + if (verbosity() > 0) { + std::cout << " Virtual Memory API not supported on this system - Skipping." << std::endl; + std::cout << kSubTestSeparator << std::endl; + } + return; + } + + ASSERT_SUCCESS(rocrtst::GetAgentPools(&agent_pools)); + + auto pool_idx = 0; + for (auto a : agent_pools) { + for (auto p : a->pools) { + TestCreateDestroy(a->agent, p); + } + } + + if (verbosity() > 0) { + std::cout << " Subtest finished" << std::endl; + std::cout << kSubTestSeparator << std::endl; + } +} + +void VirtMemoryTestBasic::TestRefCount(hsa_agent_t agent, hsa_amd_memory_pool_t pool) { + rocrtst::pool_info_t pool_i; + hsa_device_type_t ag_type; + char ag_name[64]; + void* addrRangeUnmapped; + hsa_status_t err; + void* addrRange; + + ASSERT_SUCCESS(hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, ag_name)); + ASSERT_SUCCESS(hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &ag_type)); + ASSERT_SUCCESS(rocrtst::AcquirePoolInfo(pool, &pool_i)); + + if (ag_type != HSA_DEVICE_TYPE_GPU || !pool_i.alloc_allowed) return; + + size_t granule_size = pool_i.alloc_granule; + + ASSERT_SUCCESS(hsa_amd_vmem_address_reserve(&addrRange, 10 * granule_size, 0, 0)); + + hsa_amd_vmem_alloc_handle_t mem_handleA1; + ASSERT_SUCCESS( + hsa_amd_vmem_handle_create(pool, 10 * granule_size, MEMORY_TYPE_NONE, 0, &mem_handleA1)); + ASSERT_SUCCESS(hsa_amd_vmem_map(addrRange, 10 * granule_size, 0, mem_handleA1, 0)); + + /* Allocate duplicate handle */ + hsa_amd_vmem_alloc_handle_t mem_handleA1Dup; + ASSERT_SUCCESS(hsa_amd_vmem_retain_alloc_handle(&mem_handleA1Dup, addrRange)); + + /* Try to unmap with incorrect size */ + err = hsa_amd_vmem_unmap(addrRange, 5 * granule_size); + ASSERT_NE(err, HSA_STATUS_SUCCESS); + + ASSERT_SUCCESS(hsa_amd_vmem_handle_release(mem_handleA1)); + + /* Try to release duplicate handle twice - second time should fail */ + ASSERT_SUCCESS(hsa_amd_vmem_handle_release(mem_handleA1Dup)); + + /* Already released so should fail*/ + err = hsa_amd_vmem_handle_release(mem_handleA1Dup); + ASSERT_NE(err, HSA_STATUS_SUCCESS); + + /* Unmap with correct size - un-mapping after releasing the handle is valid */ + ASSERT_SUCCESS(hsa_amd_vmem_unmap(addrRange, 10 * granule_size)); + + /* Try to free with incorrect size */ + err = hsa_amd_vmem_address_free(addrRange, 5 * granule_size); + ASSERT_NE(err, HSA_STATUS_SUCCESS); + + /* Free with correct size */ + ASSERT_SUCCESS(hsa_amd_vmem_address_free(addrRange, 10 * granule_size)); +} + +void VirtMemoryTestBasic::TestRefCount(void) { + hsa_status_t err; + std::vector> agent_pools; + + if (verbosity() > 0) { + PrintMemorySubtestHeader("Reference Count Test"); + } + bool supp = false; + ASSERT_SUCCESS(hsa_system_get_info(HSA_AMD_SYSTEM_INFO_VIRTUAL_MEM_API_SUPPORTED, (void*)&supp)); + if (!supp) { + if (verbosity() > 0) { + std::cout << " Virtual Memory API not supported on this system - Skipping." << std::endl; + std::cout << kSubTestSeparator << std::endl; + } + return; + } + ASSERT_SUCCESS(rocrtst::GetAgentPools(&agent_pools)); + + auto pool_idx = 0; + for (auto a : agent_pools) { + for (auto p : a->pools) TestRefCount(a->agent, p); + } + + if (verbosity() > 0) { + std::cout << " Subtest finished" << std::endl; + std::cout << kSubTestSeparator << std::endl; + } +} + +void VirtMemoryTestBasic::TestPartialMapping(hsa_agent_t agent, hsa_amd_memory_pool_t pool) { + rocrtst::pool_info_t pool_i; + hsa_device_type_t ag_type; + char ag_name[64]; + void* addrRangeUnmapped; + hsa_status_t err; + void* addrRange; + + ASSERT_SUCCESS(hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &ag_type)); + + ASSERT_SUCCESS(rocrtst::AcquirePoolInfo(pool, &pool_i)); + + if (ag_type != HSA_DEVICE_TYPE_GPU || !pool_i.alloc_allowed) return; + + size_t granule_size = pool_i.alloc_granule; + + /************************************************************************************************ + Map partial chunks within the address range and confirm what overlaps fail. + Units below are in multiples of granule_size. + + ------------------------------------------------------------------ + | 0 | 1 | 2 | 3 | 4 | 5 | 6 | 7 | 8 | 9 | 10 | 11 | 12 | 13 | 14 | + ------------------------------------------------------------------ + Step 1: A A A A A A + Step 2: B B B + Step 3: B + Step 4: B B B + + ***********************************************************************************************/ + + ASSERT_SUCCESS(hsa_amd_vmem_address_reserve(&addrRange, 15 * granule_size, 0, 0)); + + hsa_amd_vmem_alloc_handle_t mem_handleA; + + // Step 1 + ASSERT_SUCCESS( + hsa_amd_vmem_handle_create(pool, 8 * granule_size, MEMORY_TYPE_NONE, 0, &mem_handleA)); + + ASSERT_SUCCESS(hsa_amd_vmem_map((void*)((uint64_t)addrRange + (2 * granule_size)), + 6 * granule_size, 0, mem_handleA, 0)); + + // Step 2 + hsa_amd_vmem_alloc_handle_t mem_handleB; + ASSERT_SUCCESS( + hsa_amd_vmem_handle_create(pool, 8 * granule_size, MEMORY_TYPE_NONE, 0, &mem_handleB)); + + ASSERT_SUCCESS(hsa_amd_vmem_map((void*)((uint64_t)addrRange + (11 * granule_size)), + 3 * granule_size, 0, mem_handleB, 0)); + + // Step 3 + // Should fail as this is exceeding size of address range + err = hsa_amd_vmem_map((void*)((uint64_t)addrRange + (14 * granule_size)), + 2 * granule_size, 0, mem_handleB, 0); + ASSERT_NE(err, HSA_STATUS_SUCCESS); + + ASSERT_SUCCESS(hsa_amd_vmem_map((void*)((uint64_t)addrRange + (14 * granule_size)), + 1 * granule_size, 0, mem_handleB, 0)); + + // Step 4 + // Should fail as this is overlapping with AddressRange[11] already mapped + err = hsa_amd_vmem_map((void*)((uint64_t)addrRange + (8 * granule_size)), + 4 * granule_size, 0, mem_handleB, 0); + ASSERT_NE(err, HSA_STATUS_SUCCESS); + + ASSERT_SUCCESS(hsa_amd_vmem_map((void*)((uint64_t)addrRange + (8 * granule_size)), + 3 * granule_size, 0, mem_handleB, 0)); + + // Done, unmap all + ASSERT_SUCCESS( + hsa_amd_vmem_unmap((void*)((uint64_t)addrRange + (2 * granule_size)), 6 * granule_size)); + ASSERT_SUCCESS( + hsa_amd_vmem_unmap((void*)((uint64_t)addrRange + (8 * granule_size)), 3 * granule_size)); + ASSERT_SUCCESS( + hsa_amd_vmem_unmap((void*)((uint64_t)addrRange + (11 * granule_size)), 3 * granule_size)); + ASSERT_SUCCESS( + hsa_amd_vmem_unmap((void*)((uint64_t)addrRange + (14 * granule_size)), 1 * granule_size)); + ASSERT_SUCCESS(hsa_amd_vmem_address_free(addrRange, 15 * granule_size)); +} + +void VirtMemoryTestBasic::TestPartialMapping(void) { + hsa_status_t err; + std::vector> agent_pools; + + if (verbosity() > 0) { + PrintMemorySubtestHeader("Partial Mapping Test"); + } + + bool supp = false; + ASSERT_SUCCESS(hsa_system_get_info(HSA_AMD_SYSTEM_INFO_VIRTUAL_MEM_API_SUPPORTED, (void*)&supp)); + if (!supp) { + if (verbosity() > 0) { + std::cout << " Virtual Memory API not supported on this system - Skipping." << std::endl; + std::cout << kSubTestSeparator << std::endl; + } + return; + } + + ASSERT_SUCCESS(rocrtst::GetAgentPools(&agent_pools)); + + auto pool_idx = 0; + for (auto a : agent_pools) { + for (auto p : a->pools) TestPartialMapping(a->agent, p); + } + + if (verbosity() > 0) { + std::cout << " Subtest finished" << std::endl; + std::cout << kSubTestSeparator << std::endl; + } +} + +typedef struct __attribute__((aligned(16))) args_t { + int* a; + int* b; + int* c; +} args; + +args* kernArgsVirt = NULL; + +// Test to check CPU can read & write to GPU memory +void VirtMemoryTestBasic::CPUAccessToGPUMemoryTest(hsa_agent_t cpuAgent, hsa_agent_t gpuAgent, + hsa_amd_memory_pool_t device_pool) { + hsa_status_t err; + + rocrtst::pool_info_t pool_i; + ASSERT_SUCCESS(rocrtst::AcquirePoolInfo(device_pool, &pool_i)); + + if (!(pool_i.segment == HSA_AMD_SEGMENT_GLOBAL && + pool_i.global_flag == HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED)) + return; + + hsa_amd_memory_pool_access_t access; + hsa_amd_agent_memory_pool_get_info(cpuAgent, device_pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, + &access); + if (access == HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED) { + if (verbosity() > 0) { + std::cout << " Test not applicable as system is not large bar - Skipping." << std::endl; + std::cout << kSubTestSeparator << std::endl; + } + return; + } + if (!pool_i.alloc_allowed || pool_i.alloc_granule == 0 || pool_i.alloc_alignment == 0) { + if (verbosity() > 0) { + std::cout << " Test not applicable. Skipping." << std::endl; + std::cout << kSubTestSeparator << std::endl; + } + return; + } + + auto max_alloc_size = pool_i.alloc_granule * 100; + unsigned int max_element = max_alloc_size / sizeof(unsigned int); + unsigned int* dev_data = NULL; + unsigned int* host_data = NULL; + host_data = (unsigned int*)malloc(max_alloc_size); + + for (unsigned int i = 0; i < max_element; ++i) { + host_data[i] = i; + } + + hsa_amd_memory_access_desc_t permsAccess[] = {{HSA_ACCESS_PERMISSION_RW, cpuAgent}, + {HSA_ACCESS_PERMISSION_RW, gpuAgent}}; + + hsa_amd_vmem_alloc_handle_t mem_handle_host, mem_handle_dev; + ASSERT_SUCCESS( + hsa_amd_vmem_address_reserve(reinterpret_cast(&dev_data), max_alloc_size, 0, 0)); + ASSERT_SUCCESS(hsa_amd_vmem_handle_create(device_pool, max_alloc_size, MEMORY_TYPE_NONE, 0, + &mem_handle_dev)); + ASSERT_SUCCESS( + hsa_amd_vmem_map(reinterpret_cast(dev_data), max_alloc_size, 0, mem_handle_dev, 0)); + + // Give device access to host data + ASSERT_SUCCESS(hsa_amd_vmem_set_access(dev_data, max_alloc_size, permsAccess, 2)); + + // Verify CPU can read & write to GPU memory + std::cout << " Verify CPU can read & write to GPU memory" << std::endl; + for (unsigned int i = 0; i < max_element; ++i) { + dev_data[i] = i; // Write to gpu memory directly + } + + for (unsigned int i = 0; i < max_element; ++i) { + if (host_data[i] != dev_data[i]) { // Reading GPU memory + fprintf(stdout, + " Values not mathing !! host_data[%d]:%d ," + "dev_data[%d]\n", + host_data[i], i, dev_data[i]); + } + } + std::cout << " CPU have read & write to GPU memory successfully" << std::endl; + + ASSERT_SUCCESS(hsa_amd_vmem_unmap(dev_data, max_alloc_size)); + ASSERT_SUCCESS(hsa_amd_vmem_handle_release(mem_handle_dev)); + ASSERT_SUCCESS(hsa_amd_vmem_address_free(reinterpret_cast(dev_data), max_alloc_size)); + free(host_data); +} + +void VirtMemoryTestBasic::CPUAccessToGPUMemoryTest(void) { + hsa_status_t err; + // find all cpu agents + std::vector cpus; + ASSERT_SUCCESS(hsa_iterate_agents(rocrtst::IterateCPUAgents, &cpus)); + + // find all gpu agents + std::vector gpus; + ASSERT_SUCCESS(hsa_iterate_agents(rocrtst::IterateGPUAgents, &gpus)); + + if (verbosity() > 0) PrintMemorySubtestHeader("CPU To GPU Access test"); + + bool supp = false; + ASSERT_SUCCESS(hsa_system_get_info(HSA_AMD_SYSTEM_INFO_VIRTUAL_MEM_API_SUPPORTED, (void*)&supp)); + if (!supp) { + if (verbosity() > 0) { + std::cout << " Virtual Memory API not supported on this system - Skipping." << std::endl; + std::cout << kSubTestSeparator << std::endl; + } + return; + } + + for (unsigned int i = 0; i < gpus.size(); ++i) { + hsa_amd_memory_pool_t gpu_pool; + memset(&gpu_pool, 0, sizeof(gpu_pool)); + ASSERT_SUCCESS( + hsa_amd_agent_iterate_memory_pools(gpus[i], rocrtst::GetGlobalMemoryPool, &gpu_pool)); + if (gpu_pool.handle == 0) { + std::cout << " No global mempool in gpu agent" << std::endl; + return; + } + CPUAccessToGPUMemoryTest(cpus[0], gpus[i], gpu_pool); + } + if (verbosity() > 0) { + std::cout << " Subtest finished" << std::endl; + std::cout << kSubTestSeparator << std::endl; + } +} + +// Test to check GPU can read & write to CPU memory +void VirtMemoryTestBasic::GPUAccessToCPUMemoryTest(hsa_agent_t cpuAgent, hsa_agent_t gpuAgent, + hsa_amd_memory_pool_t device_pool) { + rocrtst::pool_info_t pool_i; + hsa_device_type_t ag_type; + char ag_name[64]; + hsa_status_t err; + + ASSERT_SUCCESS(rocrtst::AcquirePoolInfo(device_pool, &pool_i)); + + if (!pool_i.alloc_allowed || pool_i.segment != HSA_AMD_SEGMENT_GLOBAL || + pool_i.global_flag != HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED) + return; + + hsa_amd_memory_pool_access_t access; + ASSERT_SUCCESS(hsa_amd_agent_memory_pool_get_info( + cpuAgent, device_pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access)); + + if (access == HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED) { + if (verbosity() > 0) { + std::cout << " Test not applicable as system is not large bar - Skipping." << std::endl; + std::cout << kSubTestSeparator << std::endl; + return; + } + } + + hsa_queue_t* queue = NULL; // command queue + hsa_signal_t signal = {0}; // completion signal + + // size_t granule_size = pool_i.alloc_granule; + size_t granule_size = 4096; // TODO: Fixme + size_t alloc_size = granule_size * 100; + static const int kMemoryAllocSize = 1024; + // static const int kMemoryAllocSize = 4096; + unsigned int max_element = alloc_size / sizeof(unsigned int); + + // get queue size + uint32_t queue_size = 0; + ASSERT_SUCCESS(hsa_agent_get_info(gpuAgent, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_size)); + + // create queue + ASSERT_SUCCESS( + hsa_queue_create(gpuAgent, queue_size, HSA_QUEUE_TYPE_MULTI, NULL, NULL, 0, 0, &queue)); + + // Find a memory pool that supports kernel arguments. + hsa_amd_memory_pool_t kernarg_pool; + ASSERT_SUCCESS( + hsa_amd_agent_iterate_memory_pools(cpuAgent, rocrtst::GetKernArgMemoryPool, &kernarg_pool)); + + // Get System Memory Pool on the cpuAgent to allocate host side buffers + hsa_amd_memory_pool_t global_pool; + ASSERT_SUCCESS( + hsa_amd_agent_iterate_memory_pools(cpuAgent, rocrtst::GetGlobalMemoryPool, &global_pool)); + + struct host_data_t { + int data[kMemoryAllocSize * 4]; + int dup_data[kMemoryAllocSize * 4]; + int result[kMemoryAllocSize * 4]; + }; + + struct dev_data_t { + int result[kMemoryAllocSize * 4]; + }; + + + struct host_data_t* host_data; + struct dev_data_t* dev_data; + + ASSERT_SUCCESS(hsa_amd_memory_pool_allocate(global_pool, sizeof(*host_data), 0, + reinterpret_cast(&host_data))); + + // Allow gpuAgent access to all allocated system memory. + ASSERT_SUCCESS(hsa_amd_agents_allow_access(1, &gpuAgent, NULL, host_data)); + ASSERT_SUCCESS(hsa_amd_vmem_address_reserve((void**)&dev_data, sizeof(*dev_data), 0, 0)); + + hsa_amd_vmem_alloc_handle_t mem_handle; + + ASSERT_SUCCESS( + hsa_amd_vmem_handle_create(device_pool, sizeof(*dev_data), MEMORY_TYPE_NONE, 0, &mem_handle)); + ASSERT_SUCCESS(hsa_amd_vmem_map(dev_data, sizeof(*dev_data), 0, mem_handle, 0)); + + // Give host and device access to device data + hsa_amd_memory_access_desc_t permsAccess[] = {{HSA_ACCESS_PERMISSION_RW, gpuAgent}, + {HSA_ACCESS_PERMISSION_RW, cpuAgent}}; + + ASSERT_SUCCESS(hsa_amd_vmem_set_access(dev_data, sizeof(*dev_data), permsAccess, 2)); + + // Allocate the kernel argument buffer from the kernarg_pool. + ASSERT_SUCCESS(hsa_amd_memory_pool_allocate(kernarg_pool, sizeof(args_t), 0, + reinterpret_cast(&kernArgsVirt))); + + // initialize the host buffers + for (int i = 0; i < kMemoryAllocSize; ++i) { + unsigned int seed = time(NULL); + host_data->data[i] = 1 + rand_r(&seed) % 1; + host_data->dup_data[i] = host_data->data[i]; + } + + memset(host_data->result, 0, sizeof(host_data->result)); + memset(dev_data->result, 0, sizeof(dev_data->result)); + + ASSERT_SUCCESS(hsa_amd_agents_allow_access(1, &gpuAgent, NULL, kernArgsVirt)); + + kernArgsVirt->a = host_data->data; + kernArgsVirt->b = host_data->result; // system memory passed to gpu for write + kernArgsVirt->c = dev_data->result; // gpu memory to verify that gpu read system data + + // Create the executable, get symbol by name and load the code object + set_kernel_file_name("gpuReadWrite_kernels.hsaco"); + set_kernel_name("gpuReadWrite"); + ASSERT_SUCCESS(rocrtst::LoadKernelFromObjFile(this, &gpuAgent)); + + // Fill the dispatch packet with + // workgroup_size, grid_size, kernelArgs and completion signal + // Put it on the queue and launch the kernel by ringing the doorbell + + // create completion signal + ASSERT_SUCCESS(hsa_signal_create(1, 0, NULL, &signal)); + + // create aql packet + hsa_kernel_dispatch_packet_t aql; + memset(&aql, 0, sizeof(aql)); + + // initialize aql packet + aql.workgroup_size_x = 256; + aql.workgroup_size_y = 1; + aql.workgroup_size_z = 1; + aql.grid_size_x = kMemoryAllocSize; + aql.grid_size_y = 1; + aql.grid_size_z = 1; + aql.private_segment_size = 0; + aql.group_segment_size = 0; + aql.kernel_object = kernel_object(); // kernel_code; + aql.kernarg_address = kernArgsVirt; + aql.completion_signal = signal; + + // const uint32_t queue_size = queue->size; + const uint32_t queue_mask = queue->size - 1; + + // write to command queue + uint64_t index = hsa_queue_load_write_index_relaxed(queue); + hsa_queue_store_write_index_relaxed(queue, index + 1); + + rocrtst::WriteAQLToQueueLoc(queue, index, &aql); + + hsa_kernel_dispatch_packet_t* q_base_addr = + reinterpret_cast(queue->base_address); + rocrtst::AtomicSetPacketHeader( + (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | + (1 << HSA_PACKET_HEADER_BARRIER) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE), + (1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS), + reinterpret_cast(&q_base_addr[index & queue_mask])); + + // ringdoor bell + hsa_signal_store_relaxed(queue->doorbell_signal, index); + // wait for the signal and reset it for future use + while (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, 1, (uint64_t)-1, + HSA_WAIT_STATE_ACTIVE)) { + } + hsa_signal_store_relaxed(signal, 1); + + // compare device and host side results + if (verbosity() > 0) { + std::cout << " Check GPU has read the system memory" << std::endl; + } + for (int i = 0; i < kMemoryAllocSize; ++i) { + // printf("Verifying data at index[%d]\n", i); + ASSERT_EQ(dev_data->result[i], host_data->dup_data[i]); + } + + if (verbosity() > 0) { + std::cout << " GPU has read the system memory successfully" << std::endl; + std::cout << " Check GPU has written to system memory" << std::endl; + } + for (int i = 0; i < kMemoryAllocSize; ++i) { + ASSERT_EQ(host_data->result[i], i); + } + + if (verbosity() > 0) { + std::cout << " GPU has written to system memory successfully" << std::endl; + } + + ASSERT_SUCCESS(hsa_amd_vmem_unmap(dev_data, sizeof(*dev_data))); + ASSERT_SUCCESS(hsa_amd_vmem_handle_release(mem_handle)); + + if (dev_data) ASSERT_SUCCESS(hsa_amd_vmem_address_free(dev_data, sizeof(*dev_data))); + + if (host_data) hsa_memory_free(host_data); + if (kernArgsVirt) { + hsa_memory_free(kernArgsVirt); + } + if (signal.handle) { + hsa_signal_destroy(signal); + } + if (queue) { + hsa_queue_destroy(queue); + } +} + +void VirtMemoryTestBasic::GPUAccessToCPUMemoryTest(void) { + hsa_status_t err; + // find all cpu agents + std::vector cpus; + ASSERT_SUCCESS(hsa_iterate_agents(rocrtst::IterateCPUAgents, &cpus)); + + // find all gpu agents + std::vector gpus; + ASSERT_SUCCESS(hsa_iterate_agents(rocrtst::IterateGPUAgents, &gpus)); + + if (verbosity() > 0) PrintMemorySubtestHeader("CPU To GPU Access test"); + + bool supp = false; + ASSERT_SUCCESS(hsa_system_get_info(HSA_AMD_SYSTEM_INFO_VIRTUAL_MEM_API_SUPPORTED, (void*)&supp)); + if (!supp) { + if (verbosity() > 0) { + std::cout << " Virtual Memory API not supported on this system - Skipping." << std::endl; + std::cout << kSubTestSeparator << std::endl; + } + return; + } + + for (unsigned int i = 0; i < gpus.size(); ++i) { + hsa_amd_memory_pool_t gpu_pool; + memset(&gpu_pool, 0, sizeof(gpu_pool)); + ASSERT_SUCCESS( + hsa_amd_agent_iterate_memory_pools(gpus[i], rocrtst::GetGlobalMemoryPool, &gpu_pool)); + if (gpu_pool.handle == 0) { + std::cout << "no global mempool in GPU agent" << std::endl; + return; + } + GPUAccessToCPUMemoryTest(cpus[0], gpus[i], gpu_pool); + } + if (verbosity() > 0) { + std::cout << " Subtest finished" << std::endl; + std::cout << kSubTestSeparator << std::endl; + } +} + +void VirtMemoryTestBasic::SetUp(void) { + hsa_status_t err; + + TestBase::SetUp(); + + ASSERT_SUCCESS(rocrtst::SetDefaultAgents(this)); + ASSERT_SUCCESS(rocrtst::SetPoolsTypical(this)); + + return; +} + +void VirtMemoryTestBasic::Run(void) { + // Compare required profile for this test case with what we're actually + // running on + if (!rocrtst::CheckProfile(this)) { + return; + } + + TestBase::Run(); +} + +void VirtMemoryTestBasic::DisplayTestInfo(void) { TestBase::DisplayTestInfo(); } + +void VirtMemoryTestBasic::DisplayResults(void) const { + // Compare required profile for this test case with what we're actually + // running on + if (!rocrtst::CheckProfile(this)) { + return; + } + + return; +} + +void VirtMemoryTestBasic::Close() { + // This will close handles opened within rocrtst utility calls and call + // hsa_shut_down(), so it should be done after other hsa cleanup + TestBase::Close(); +} diff --git a/projects/rocr-runtime/rocrtst/suites/functional/virtual_memory.h b/projects/rocr-runtime/rocrtst/suites/functional/virtual_memory.h new file mode 100644 index 0000000000..2b9b243e3e --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/functional/virtual_memory.h @@ -0,0 +1,93 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2022, Advanced Micro Devices, Inc. + * All rights reserved. + * + * Developed by: + * + * AMD Research and AMD ROC Software Development + * + * Advanced Micro Devices, Inc. + * + * www.amd.com + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal with the Software without restriction, including without limitation + * the rights to use, copy, modify, merge, publish, distribute, sublicense, + * and/or sell copies of the Software, and to permit persons to whom the + * Software is furnished to do so, subject to the following conditions: + * + * - Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimers. + * - Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimers in + * the documentation and/or other materials provided with the distribution. + * - Neither the names of , + * nor the names of its contributors may be used to endorse or promote + * products derived from this Software without specific prior written + * permission. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL + * THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR + * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, + * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER + * DEALINGS WITH THE SOFTWARE. + * + */ +#ifndef ROCRTST_SUITES_FUNCTIONAL_VIRTUAL_MEMORY_H_ +#define ROCRTST_SUITES_FUNCTIONAL_VIRTUAL_MEMORY_H_ + +#include + +#include "common/base_rocr.h" +#include "hsa/hsa.h" +#include "suites/test_common/test_base.h" + +class VirtMemoryTestBasic : public TestBase { + public: + VirtMemoryTestBasic(); + + // @Brief: Destructor for test case of VirtMemoryTestBasic + virtual ~VirtMemoryTestBasic(); + + // @Brief: Setup the environment for measurement + virtual void SetUp(); + + // @Brief: Core measurement execution + virtual void Run(); + + // @Brief: Clean up and retrive the resource + virtual void Close(); + + // @Brief: Display results + virtual void DisplayResults() const; + + // @Brief: Display information about what this test does + virtual void DisplayTestInfo(void); + + void TestCreateDestroy(void); + void TestRefCount(void); + void TestPartialMapping(void); + void GPUAccessToCPUMemoryTest(void); + void CPUAccessToGPUMemoryTest(void); + + private: + void TestCreateDestroy(hsa_agent_t agent, hsa_amd_memory_pool_t pool); + void TestRefCount(hsa_agent_t agent, hsa_amd_memory_pool_t pool); + void TestPartialMapping(hsa_agent_t agent, hsa_amd_memory_pool_t pool); + void GPUAccessToCPUMemoryTest(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent, + hsa_amd_memory_pool_t pool); + void CPUAccessToGPUMemoryTest(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent, + hsa_amd_memory_pool_t pool); +}; + + +#endif // ROCRTST_SUITES_FUNCTIONAL_VIRTUAL_MEMORY_H_ diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/main.cc b/projects/rocr-runtime/rocrtst/suites/test_common/main.cc index 848f6f5927..5a3b1841fa 100755 --- a/projects/rocr-runtime/rocrtst/suites/test_common/main.cc +++ b/projects/rocr-runtime/rocrtst/suites/test_common/main.cc @@ -57,6 +57,7 @@ #include "suites/functional/memory_atomics.h" #include "suites/functional/memory_allocation.h" #include "suites/functional/deallocation_notifier.h" +#include "suites/functional/virtual_memory.h" #include "suites/performance/dispatch_time.h" #include "suites/performance/memory_async_copy.h" #include "suites/performance/memory_async_copy_numa.h" @@ -373,6 +374,18 @@ TEST(rocrtstFunc, AgentProp_UUID) { RunCustomTestEpilog(&propTest); } +TEST(rocrtstFunc, VirtMemory_Basic_Test) { + VirtMemoryTestBasic vmt; + + RunCustomTestProlog(&vmt); + vmt.TestCreateDestroy(); + vmt.TestRefCount(); + vmt.TestPartialMapping(); + vmt.CPUAccessToGPUMemoryTest(); + vmt.GPUAccessToCPUMemoryTest(); + RunCustomTestEpilog(&vmt); +} + TEST(rocrtstNeg, Memory_Negative_Tests) { MemoryAllocateNegativeTest mt; RunCustomTestProlog(&mt);