AI/rocm-systems
2
0
Çatalla 0
Kod Konular Değişiklik İstekleri İşlemler Paketler Projeler Sürüm Viki Aktivite

Modified memory atomics on non Large bar system and pool info test memcmp issue fix

Kaynağa Gözat 
Change-Id: I951fdb6c91508f43b1c51f7eb92870543fc58e53
Bu işleme şunda yer alıyor:
rohit pathania
2018-05-31 15:26:01 +05:30
ebeveyn 12a81ae96f
işleme d8e47ba115
4 değiştirilmiş dosya ile 322 ekleme ve 248 silme
Yama Dosyasını İndir Diff Dosyasını İndir Tüm dosyaları genişlet Tüm dosyaları daralt
rocrtst/common/common.h
+12 -1
Dosyayı Görüntüle
@@ -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;
rocrtst/suites/functional/memory_atomics.cc
+308 -232
Dosyayı Görüntüle
@@ -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<void **>(&oldValues));
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
err = hsa_amd_memory_pool_allocate(global_pool,
kMemoryAllocSize, 0,
reinterpret_cast<void **>(&refSysdata));
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
err = hsa_amd_memory_pool_allocate(global_pool,
kMemoryAllocSize, 0,
reinterpret_cast<void **>(&oldrefdata));
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
err = hsa_amd_memory_pool_allocate(global_pool,
kMemoryAllocSize, 0,
reinterpret_cast<void **>(&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<void **>(&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<void **>(&oldValues));
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
err = hsa_amd_memory_pool_allocate(global_pool,
kMemoryAllocSize, 0,
reinterpret_cast<void **>(&refSysdata));
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
err = hsa_amd_memory_pool_allocate(global_pool,
kMemoryAllocSize, 0,
reinterpret_cast<void **>(&oldrefdata));
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
err = hsa_amd_memory_pool_allocate(global_pool,
kMemoryAllocSize, 0,
reinterpret_cast<void **>(&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<void **>(&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<void **>(&gpuRefData));
reinterpret_cast<void **>(&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, &copy_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<void **>(&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<void **>(&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<void **>(&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" <<std::endl;
}
} else if (testtype_ == SUB) {
set_kernel_name("test_atomic_sub");
// set the expected data result set from kernel
for (int i = 0; i < arraySize; ++i) {
expecteddata[i] = oldrefdata[i] - kValue;
}
err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
// Fill up the kernel packet except header
err = rocrtst::InitializeAQLPacket(this, &aql());
ASSERT_EQ(HSA_STATUS_SUCCESS, err);
aql().workgroup_size_x = 256;
aql().workgroup_size_y = 1;
aql().workgroup_size_z = 1;
aql().grid_size_x = kMemoryAllocSize;
aql().kernarg_address = kernArguments;
aql().kernel_object = kernel_object();
const uint32_t queue_mask = queue->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<hsa_kernel_dispatch_packet_t*>
(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" <<std::endl;
}
return;
}
err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
// Fill up the kernel packet except header
err = rocrtst::InitializeAQLPacket(this, &aql());
ASSERT_EQ(HSA_STATUS_SUCCESS, err);
aql().workgroup_size_x = 256;
aql().workgroup_size_y = 1;
aql().workgroup_size_z = 1;
aql().grid_size_x = arraySize;
aql().kernarg_address = kernArguments;
aql().kernel_object = kernel_object();
const uint32_t queue_mask = queue->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<hsa_kernel_dispatch_packet_t*>
(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
rocrtst/suites/stress/memory_concurrent_tests.cc
+1 -14
Dosyayı Görüntüle
@@ -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
rocrtst/suites/test_common/main.cc
+1 -1
Dosyayı Görüntüle
@@ -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();