From d8e47ba115a6c54c577e48089d79463e9a03fa1b Mon Sep 17 00:00:00 2001 From: rohit pathania Date: Thu, 31 May 2018 15:26:01 +0530 Subject: [PATCH] Modified memory atomics on non Large bar system and pool info test memcmp issue fix Change-Id: I951fdb6c91508f43b1c51f7eb92870543fc58e53 --- rocrtst/common/common.h | 13 +- rocrtst/suites/functional/memory_atomics.cc | 540 ++++++++++-------- .../suites/stress/memory_concurrent_tests.cc | 15 +- rocrtst/suites/test_common/main.cc | 2 +- 4 files changed, 322 insertions(+), 248 deletions(-) diff --git a/rocrtst/common/common.h b/rocrtst/common/common.h index 475ab2113b..378156a989 100755 --- a/rocrtst/common/common.h +++ b/rocrtst/common/common.h @@ -80,7 +80,7 @@ namespace rocrtst { // This structure holds memory pool information acquired through hsa info // related calls, and is later used for reference when displaying the // information. -typedef struct { +typedef struct pool_info_t_ { uint32_t segment; size_t size; bool alloc_allowed; @@ -88,6 +88,17 @@ typedef struct { size_t alloc_alignment; bool accessible_by_all; uint32_t global_flag; + inline bool operator==(const pool_info_t_ &a) { + if (a.segment == segment && a.size == size + && a.alloc_allowed == alloc_allowed + && a.alloc_granule == alloc_granule + && a.alloc_alignment == alloc_alignment + && a.accessible_by_all == accessible_by_all + && a.global_flag == global_flag ) + return true; + else + return false; + } } pool_info_t; diff --git a/rocrtst/suites/functional/memory_atomics.cc b/rocrtst/suites/functional/memory_atomics.cc index 943766bd74..0fc498f182 100755 --- a/rocrtst/suites/functional/memory_atomics.cc +++ b/rocrtst/suites/functional/memory_atomics.cc @@ -212,7 +212,7 @@ typedef struct __attribute__ ((aligned(16))) args_t { static const char kSubTestSeparator[] = " **************************"; -static const int kMemoryAllocSize = 10; +static const int kMemoryAllocSize = 4096; void MemoryAtomic::MemoryAtomicTest(hsa_agent_t cpuAgent, hsa_agent_t gpuAgent) { @@ -229,257 +229,333 @@ void MemoryAtomic::MemoryAtomicTest(hsa_agent_t cpuAgent, hsa_amd_agent_memory_pool_get_info(cpuAgent, gpu_pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access); + // hsa objects + hsa_queue_t *queue = NULL; // command queue + // get queue size + uint32_t queue_size = 0; + err = hsa_agent_get_info(gpuAgent, + HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_size); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + // create queue + err = hsa_queue_create(gpuAgent, + queue_size, HSA_QUEUE_TYPE_MULTI, + NULL, NULL, 0, 0, &queue); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + // Find a memory pool that supports kernel arguments. + hsa_amd_memory_pool_t kernarg_pool; + err = hsa_amd_agent_iterate_memory_pools(cpuAgent, + rocrtst::GetKernArgMemoryPool, + &kernarg_pool); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + // Allocate the host side buffers + // (refSysdata,oldValues,oldrefdata,kernArg) on system memory + + // this is ref sys data on which atomics operation need to done + int *refSysdata = NULL; + // This is oldrefdata which will be required to compare the returned old values after atomics operation + int *oldrefdata = NULL; + // This is returned old values + int *oldValues = NULL; + // This is expected data set + int *expecteddata = NULL; + // Array size for the data + int arraySize = kMemoryAllocSize/sizeof(int); + + // Get System Memory Pool on the cpuAgent to allocate host side buffers + hsa_amd_memory_pool_t global_pool; + err = hsa_amd_agent_iterate_memory_pools(cpuAgent, + rocrtst::GetGlobalMemoryPool, + &global_pool); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + err = hsa_amd_memory_pool_allocate(global_pool, + kMemoryAllocSize, 0, + reinterpret_cast(&oldValues)); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + err = hsa_amd_memory_pool_allocate(global_pool, + kMemoryAllocSize, 0, + reinterpret_cast(&refSysdata)); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + err = hsa_amd_memory_pool_allocate(global_pool, + kMemoryAllocSize, 0, + reinterpret_cast(&oldrefdata)); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + err = hsa_amd_memory_pool_allocate(global_pool, + kMemoryAllocSize, 0, + reinterpret_cast(&expecteddata)); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + + // Allocate the kernel argument buffer from the kernarg_pool. + args *kernArguments = NULL; + err = hsa_amd_memory_pool_allocate(kernarg_pool, sizeof(args_t), 0, + reinterpret_cast(&kernArguments)); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + + memset(oldValues, 0, kMemoryAllocSize); + memset(expecteddata, 0, kMemoryAllocSize); + // this signal will be used for copying the data memory from To and fro from GPU + // on Non-largebar system + hsa_signal_t copy_signal; + + // for the dGPU, we have coarse grained local memory, + // so allocate memory for it on the GPU's GLOBAL segment . + + // Get local memory of GPU to allocate device side buffers on which atomics operation need to done + int *gpuRefData = NULL; + + // On non-Large bar system acess to GPU pool not allowed to directly so pinned memory + // g_gpuRefData is pointer to GPU Memory allocated on non-large bar where + // gpuRefData would be pointer to host allocated memory on non-large bar + int *g_gpuRefData = NULL; + // Pointer to the location where to store the new address + int *device_ptr = NULL; + if (access != HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED) { - // hsa objects - hsa_queue_t *queue = NULL; // command queue - // get queue size - uint32_t queue_size = 0; - err = hsa_agent_get_info(gpuAgent, - HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_size); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); - - // create queue - err = hsa_queue_create(gpuAgent, - queue_size, HSA_QUEUE_TYPE_MULTI, - NULL, NULL, 0, 0, &queue); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); - - // Find a memory pool that supports kernel arguments. - hsa_amd_memory_pool_t kernarg_pool; - err = hsa_amd_agent_iterate_memory_pools(cpuAgent, - rocrtst::GetKernArgMemoryPool, - &kernarg_pool); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); - - // Allocate the host side buffers - // (refSysdata,oldValues,oldrefdata,kernArg) on system memory - - // this is ref sys data on which atomics operation need to done - int *refSysdata = NULL; - // This is oldrefdata which will be required to compare the returned old values after atomics operation - int *oldrefdata = NULL; - // This is returned old values - int *oldValues = NULL; - // This is expected data set - int *expecteddata = NULL; - - // Get System Memory Pool on the cpuAgent to allocate host side buffers - hsa_amd_memory_pool_t global_pool; - err = hsa_amd_agent_iterate_memory_pools(cpuAgent, - rocrtst::GetGlobalMemoryPool, - &global_pool); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); - - err = hsa_amd_memory_pool_allocate(global_pool, - kMemoryAllocSize, 0, - reinterpret_cast(&oldValues)); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); - - err = hsa_amd_memory_pool_allocate(global_pool, - kMemoryAllocSize, 0, - reinterpret_cast(&refSysdata)); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); - - err = hsa_amd_memory_pool_allocate(global_pool, - kMemoryAllocSize, 0, - reinterpret_cast(&oldrefdata)); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); - - err = hsa_amd_memory_pool_allocate(global_pool, - kMemoryAllocSize, 0, - reinterpret_cast(&expecteddata)); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); - - - // Allocate the kernel argument buffer from the kernarg_pool. - args *kernArguments = NULL; - err = hsa_amd_memory_pool_allocate(kernarg_pool, sizeof(args_t), 0, - reinterpret_cast(&kernArguments)); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); - - - memset(oldValues, 0, kMemoryAllocSize * sizeof(int)); - memset(expecteddata, 0, kMemoryAllocSize * sizeof(int)); - - // for the dGPU, we have coarse grained local memory, - // so allocate memory for it on the GPU's GLOBAL segment . - - // Get local memory of GPU to allocate device side buffers on which atomics operation need to done - int *gpuRefData = NULL; err = hsa_amd_memory_pool_allocate(gpu_pool, kMemoryAllocSize, 0, - reinterpret_cast(&gpuRefData)); + reinterpret_cast(&gpuRefData)); ASSERT_EQ(err, HSA_STATUS_SUCCESS); - // Allow cpuAgent access to all allocated GPU memory. err = hsa_amd_agents_allow_access(1, &cpuAgent, NULL, gpuRefData); ASSERT_EQ(err, HSA_STATUS_SUCCESS); - memset(gpuRefData, 0, kMemoryAllocSize * sizeof(int)); - - // initialize the host buffers & gpuRefData buffer - for (int i = 0; i < kMemoryAllocSize; ++i) { - unsigned int seed = time(NULL); - refSysdata[i] = 6 + rand_r(&seed) % 1; - gpuRefData[i] = 6 + rand_r(&seed) % 1; - oldrefdata[i] = refSysdata[i]; - } + memset(gpuRefData, 0, kMemoryAllocSize); + } else { + err = hsa_signal_create(1, 0, NULL, ©_signal); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + // Alocate the System Memory and get pointer gpuRefData + err = hsa_amd_memory_pool_allocate(global_pool, kMemoryAllocSize, 0, + reinterpret_cast(&gpuRefData)); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + memset(gpuRefData, 0, kMemoryAllocSize); + // Alocate the GPU Memory and get pointer g_gpuRefData + err = hsa_amd_memory_pool_allocate(gpu_pool, kMemoryAllocSize, 0, + reinterpret_cast(&g_gpuRefData)); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + // Pinned the Host memory and get the pointer to new adress which is accesible to GPU agent + err = hsa_amd_memory_lock(gpuRefData, kMemoryAllocSize, &gpuAgent, 1, reinterpret_cast(&device_ptr)); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + } - // Allow gpuAgent access to all allocated system memory. - err = hsa_amd_agents_allow_access(1, &gpuAgent, NULL, oldValues); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); - err = hsa_amd_agents_allow_access(1, &gpuAgent, NULL, refSysdata); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); - err = hsa_amd_agents_allow_access(1, &gpuAgent, NULL, oldrefdata); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); - err = hsa_amd_agents_allow_access(1, &gpuAgent, NULL, kernArguments); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); - kernArguments->a = refSysdata; + // initialize the host buffers & gpuRefData buffer + for (int i = 0; i < arraySize; ++i) { + unsigned int seed = time(NULL); + refSysdata[i] = 6 + rand_r(&seed) % 1; + gpuRefData[i] = 6 + rand_r(&seed) % 1; + oldrefdata[i] = refSysdata[i]; + } + + // Sync the data from system memory to GPU memory on non-largebar + if (access == HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED) { + hsa_signal_store_relaxed(copy_signal, 1); + err = hsa_amd_memory_async_copy(g_gpuRefData, gpuAgent, device_ptr, + gpuAgent, kMemoryAllocSize, 0, NULL, copy_signal); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + while (hsa_signal_wait_acquire(copy_signal, HSA_SIGNAL_CONDITION_LT, 1, (uint64_t)(-1), HSA_WAIT_STATE_ACTIVE)) {} + } + + + // Allow gpuAgent access to all allocated system memory. + err = hsa_amd_agents_allow_access(1, &gpuAgent, NULL, oldValues); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + err = hsa_amd_agents_allow_access(1, &gpuAgent, NULL, refSysdata); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + err = hsa_amd_agents_allow_access(1, &gpuAgent, NULL, oldrefdata); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + err = hsa_amd_agents_allow_access(1, &gpuAgent, NULL, kernArguments); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + kernArguments->a = refSysdata; + if (access != HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED) { kernArguments->b = gpuRefData; - kernArguments->c = oldValues; - if (testtype_ != INC && testtype_ != DEC) { - kernArguments->d = kValue; + } else { + kernArguments->b = g_gpuRefData; + } + kernArguments->c = oldValues; + + if (testtype_ != INC && testtype_ != DEC) { + kernArguments->d = kValue; + } + + // Create the executable, get symbol by name and load the code object + set_kernel_file_name("atomicOperations_kernels.hsaco"); + + if (testtype_ == ADD) { + set_kernel_name("test_atomic_add"); + // set the expected data result set from kernel + for (int i = 0; i < arraySize; ++i) { + expecteddata[i] = oldrefdata[i] + kValue; } - - // Create the executable, get symbol by name and load the code object - set_kernel_file_name("atomicOperations_kernels.hsaco"); - - if (testtype_ == ADD) { - set_kernel_name("test_atomic_add"); - // set the expected data result set from kernel - for (int i = 0; i < kMemoryAllocSize; ++i) { - expecteddata[i] = oldrefdata[i] + kValue; - } - } else if (testtype_ == SUB) { - set_kernel_name("test_atomic_sub"); - // set the expected data result set from kernel - for (int i = 0; i < kMemoryAllocSize; ++i) { - expecteddata[i] = oldrefdata[i] - kValue; - } - } else if (testtype_ == AND) { - set_kernel_name("test_atomic_and"); - // set the expected data result set from kernel - for (int i = 0; i < kMemoryAllocSize; ++i) { - expecteddata[i] = oldrefdata[i] & kValue; - } - } else if (testtype_ == OR) { - set_kernel_name("test_atomic_or"); - // set the expected data result set from kernel - for (int i = 0; i < kMemoryAllocSize; ++i) { - expecteddata[i] = oldrefdata[i] | kValue; - } - } else if (testtype_ == XOR) { - set_kernel_name("test_atomic_xor"); - // set the expected data result set from kernel - for (int i = 0; i < kMemoryAllocSize; ++i) { - expecteddata[i] = oldrefdata[i] ^ kValue; - } - } else if (testtype_ == MIN) { - set_kernel_name("test_atomic_min"); - // set the expected data result set from kernel - for (int i = 0; i < kMemoryAllocSize; ++i) { - expecteddata[i] = std::min(oldrefdata[i], kValue); - } - } else if (testtype_ == MAX) { - set_kernel_name("test_atomic_max"); - // set the expected data result set from kernel - for (int i = 0; i < kMemoryAllocSize; ++i) { - expecteddata[i] = std::max(oldrefdata[i], kValue); - } - } else if (testtype_ == INC) { - set_kernel_name("test_atomic_inc"); - // set the expected data result set from kernel - for (int i = 0; i < kMemoryAllocSize; ++i) { - expecteddata[i] = oldrefdata[i] + 4; - } - } else if (testtype_ == DEC) { - set_kernel_name("test_atomic_dec"); - // set the expected data result set from kernel - for (int i = 0; i < kMemoryAllocSize; ++i) { - expecteddata[i] = oldrefdata[i] - 4; - } - } else if (testtype_ == XCHG) { - set_kernel_name("test_atomic_xchg"); - // set the expected data result set from kernel - for (int i = 0; i < kMemoryAllocSize; ++i) { - expecteddata[i] = kValue; - } - } else { - if (verbosity() > 0) { - std::cout<< "No test specified" <size - 1; - - // Load index for writing header later to command queue at same index - uint64_t index = hsa_queue_load_write_index_relaxed(queue); - - // This function simply copies the data we've collected so far into our - // local AQL packet, except the the setup and header fields. - WriteAQLPktToQueue(queue); - - - aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH; - aql().header |= HSA_FENCE_SCOPE_SYSTEM << - HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE; - aql().header |= HSA_FENCE_SCOPE_SYSTEM << - HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE; - - void* q_base = queue->base_address; - // Set the Aql packet header - AtomicSetPacketHeader(aql().header, aql().setup, - &(reinterpret_cast - (q_base))[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(aql().completion_signal, HSA_SIGNAL_CONDITION_LT, 1, - (uint64_t)-1, HSA_WAIT_STATE_ACTIVE)) { } - - hsa_signal_store_relaxed(aql().completion_signal, 1); - - // compare results with expected results - for (int i = 0; i < kMemoryAllocSize; ++i) { - ASSERT_EQ(refSysdata[i], expecteddata[i]); - ASSERT_EQ(gpuRefData[i], expecteddata[i]); - ASSERT_EQ(oldValues[i], oldrefdata[i]); + } else if (testtype_ == AND) { + set_kernel_name("test_atomic_and"); + // set the expected data result set from kernel + for (int i = 0; i < arraySize; ++i) { + expecteddata[i] = oldrefdata[i] & kValue; + } + } else if (testtype_ == OR) { + set_kernel_name("test_atomic_or"); + // set the expected data result set from kernel + for (int i = 0; i < arraySize; ++i) { + expecteddata[i] = oldrefdata[i] | kValue; + } + } else if (testtype_ == XOR) { + set_kernel_name("test_atomic_xor"); + // set the expected data result set from kernel + for (int i = 0; i < arraySize; ++i) { + expecteddata[i] = oldrefdata[i] ^ kValue; + } + } else if (testtype_ == MIN) { + set_kernel_name("test_atomic_min"); + // set the expected data result set from kernel + for (int i = 0; i < arraySize; ++i) { + expecteddata[i] = std::min(oldrefdata[i], kValue); + } + } else if (testtype_ == MAX) { + set_kernel_name("test_atomic_max"); + // set the expected data result set from kernel + for (int i = 0; i < arraySize; ++i) { + expecteddata[i] = std::max(oldrefdata[i], kValue); + } + } else if (testtype_ == INC) { + set_kernel_name("test_atomic_inc"); + // set the expected data result set from kernel + for (int i = 0; i < arraySize; ++i) { + expecteddata[i] = oldrefdata[i] + 4; + } + } else if (testtype_ == DEC) { + set_kernel_name("test_atomic_dec"); + // set the expected data result set from kernel + for (int i = 0; i < arraySize; ++i) { + expecteddata[i] = oldrefdata[i] - 4; + } + } else if (testtype_ == XCHG) { + set_kernel_name("test_atomic_xchg"); + // set the expected data result set from kernel + for (int i = 0; i < arraySize; ++i) { + expecteddata[i] = kValue; } - - if (refSysdata) { hsa_memory_free(refSysdata); } - if (oldrefdata) { hsa_memory_free(oldrefdata); } - if (oldValues) {hsa_memory_free(oldValues); } - if (gpuRefData) {hsa_memory_free(gpuRefData); } - if (kernArguments) { hsa_memory_free(kernArguments); } - if (queue) { hsa_queue_destroy(queue); } } else { if (verbosity() > 0) { - std::cout<< "Test not applicable as system is not large bar." - "Skipping."<< std::endl; - std::cout << kSubTestSeparator << std::endl; + std::cout<< "No test specified" <size - 1; + + // Load index for writing header later to command queue at same index + uint64_t index = hsa_queue_load_write_index_relaxed(queue); + + // This function simply copies the data we've collected so far into our + // local AQL packet, except the the setup and header fields. + WriteAQLPktToQueue(queue); + + + aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH; + aql().header |= HSA_FENCE_SCOPE_SYSTEM << + HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE; + aql().header |= HSA_FENCE_SCOPE_SYSTEM << + HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE; + + void* q_base = queue->base_address; + // Set the Aql packet header + AtomicSetPacketHeader(aql().header, aql().setup, + &(reinterpret_cast + (q_base))[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(aql().completion_signal, HSA_SIGNAL_CONDITION_LT, 1, + (uint64_t)-1, HSA_WAIT_STATE_ACTIVE)) { } + + hsa_signal_store_relaxed(aql().completion_signal, 1); + + // Sync the data from GPU memory to system memory on non-largebar + if (access == HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED) { + hsa_signal_store_relaxed(copy_signal, 1); + err = hsa_amd_memory_async_copy(device_ptr, gpuAgent, g_gpuRefData, + gpuAgent, kMemoryAllocSize, 0, NULL, copy_signal); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + while (hsa_signal_wait_acquire(copy_signal, HSA_SIGNAL_CONDITION_LT, 1, (uint64_t)(-1), HSA_WAIT_STATE_ACTIVE)) { } + } + + // compare results with expected results + for (int i = 0; i < arraySize; ++i) { + ASSERT_EQ(refSysdata[i], expecteddata[i]); + ASSERT_EQ(gpuRefData[i], expecteddata[i]); + ASSERT_EQ(oldValues[i], oldrefdata[i]); + } + + if (refSysdata) { + err = hsa_memory_free(refSysdata); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + } + if (oldrefdata) { + err = hsa_memory_free(oldrefdata); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + } + if (oldValues) { + err = hsa_memory_free(oldValues); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + } + if (access == HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED) { + err = hsa_amd_memory_unlock(gpuRefData); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + // Destroy the copy signal + err = hsa_signal_destroy(copy_signal); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + if (g_gpuRefData) { + err = hsa_memory_free(g_gpuRefData); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + } + } + if (gpuRefData) { + err = hsa_memory_free(gpuRefData); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + } + if (kernArguments) { + err = hsa_memory_free(kernArguments); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + } + if (queue) { + err = hsa_queue_destroy(queue); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); } } + + void MemoryAtomic::MemoryAtomicTest(void) { hsa_status_t err; // find all cpu agents diff --git a/rocrtst/suites/stress/memory_concurrent_tests.cc b/rocrtst/suites/stress/memory_concurrent_tests.cc index 1167e2f4b5..0eb213589b 100755 --- a/rocrtst/suites/stress/memory_concurrent_tests.cc +++ b/rocrtst/suites/stress/memory_concurrent_tests.cc @@ -67,17 +67,6 @@ static const uint32_t kMaxAllocSize = 1024 * 1024; -#define RET_IF_HSA_ERR(err) { \ - if ((err) != HSA_STATUS_SUCCESS) { \ - const char* msg = 0; \ - hsa_status_string(err, &msg); \ - std::cout << "hsa api call failure at line " << __LINE__ << ", file: " << \ - __FILE__ << ". Call returned " << err << std::endl; \ - std::cout << msg << std::endl; \ - return (err); \ - } \ -} - typedef struct control_block { hsa_amd_memory_pool_t* pool; @@ -134,7 +123,7 @@ static void CallbackGetPoolInfo(void* data) { err = rocrtst::AcquirePoolInfo(thread_data->pool, &info); ASSERT_EQ(HSA_STATUS_SUCCESS, err); - if (0 == memcmp(thread_data->info, &info, sizeof(rocrtst::pool_info_t))) { + if (*(thread_data->info) == info) { // The pool info is consistent with the one got from the main thread thread_data->consistency = 1; } else { @@ -533,5 +522,3 @@ void MemoryConcurrentTest::MemoryConcurrentPoolGetInfo(void) { std::cout << kSubTestSeparator << std::endl; } } - -#undef RET_IF_HSA_ERR diff --git a/rocrtst/suites/test_common/main.cc b/rocrtst/suites/test_common/main.cc index 71c473f4c1..9327a1338d 100755 --- a/rocrtst/suites/test_common/main.cc +++ b/rocrtst/suites/test_common/main.cc @@ -386,7 +386,7 @@ TEST(rocrtstStress, Memory_Concurrent_Free_Test) { RunCustomTestEpilog(&mt); } -TEST(rocrtstStress, DISABLED_Memory_Concurrent_Pool_Info_Test) { +TEST(rocrtstStress, Memory_Concurrent_Pool_Info_Test) { MemoryConcurrentTest mt(false, false, true); RunCustomTestProlog(&mt); mt.MemoryConcurrentPoolGetInfo();