AI/rocm-systems
2
0
Fork 0
Koodi Ongelmat Pull-pyynnöt Actions Paketit Projektit Julkaisut Wiki Toiminta

SWDEV-385489 - [catch2][dtest] Adding scenarios for hipEventCreateWithFlags() with hipEventDisableSystemFence and default flags (#282)

Selaa lähdekoodia 
Change-Id: Ied3dce437f34d793eb3715999125924ff74ff9b8

[ROCm/hip-tests commit: 23b02df2a1]
This commit is contained in:
ROCm CI Service Account
2023-08-14 20:53:37 +05:30
committed by GitHub
vanhempi feb18300fb
commit a4378fe258
4 muutettua tiedostoa jossa 323 lisäystä ja 11 poistoa
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/hipTestMain/config/config_amd_windows_MI2xx.json
+7
Näytä tiedosto
@@ -98,6 +98,13 @@
"Unit_hipStreamValue_Wait64_Blocking_NoMask_Nor",
"Unit_hipLaunchHostFunc_Graph",
"Unit_hipLaunchHostFunc_KernelHost",
"Unit_hipEventCreateWithFlags_DisableSystemFence_HstVisMem",
"Unit_hipEventCreateWithFlags_DefaultFlg_HstVisMem",
"Unit_hipEventCreateWithFlags_DisableSystemFence_NonCohHstMem",
"Unit_hipEventCreateWithFlags_DefaultFlg_NonCohHstMem",
"Unit_hipEventCreateWithFlags_DisableSystemFence_CohHstMem",
"Unit_hipEventCreateWithFlags_DefaultFlg_CohHstMem",
"Unit_hipStreamCreateWithPriority_MulthreadNonblockingflag",
"Unit_hipDeviceGetUuid_Positive",
"=== Below hiprtc tests are disabled temporarily, will be renabled once patches for SWDEV-395996 are merged ===",
"Unit_hiprtc_saxpy.Unit_hiprtc_saxpy",
projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json
+6
Näytä tiedosto
@@ -104,6 +104,12 @@
"Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout",
"Unit_hipFuncSetAttribute_Positive_Parameters",
"Unit_hipFuncSetAttribute_Negative_Parameters",
"Unit_hipEventCreateWithFlags_DisableSystemFence_HstVisMem",
"Unit_hipEventCreateWithFlags_DefaultFlg_HstVisMem",
"Unit_hipEventCreateWithFlags_DisableSystemFence_NonCohHstMem",
"Unit_hipEventCreateWithFlags_DefaultFlg_NonCohHstMem",
"Unit_hipEventCreateWithFlags_DisableSystemFence_CohHstMem",
"Unit_hipEventCreateWithFlags_DefaultFlg_CohHstMem",
"NOTE: The following 4 tests are disabled due to defect - EXSWHTEC-240",
"Unit_hipFuncSetCacheConfig_Negative_Not_Supported",
"Unit_hipFuncSetSharedMemConfig_Negative_Not_Supported",
projects/hip-tests/catch/include/hip_test_defgroups.hh
+7
Näytä tiedosto
@@ -107,6 +107,13 @@ THE SOFTWARE.
* @}
*/
/**
* @defgroup EventTest Event Management
* @{
* This section describes the event management types & functions of HIP runtime API.
* @}
*/
/**
* @defgroup ContextTest Context Management
* @{
projects/hip-tests/catch/unit/event/hipEventCreateWithFlags.cc
+303 -11
Näytä tiedosto
@@ -1,5 +1,5 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2022-2023 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
@@ -21,38 +21,330 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_defgroups.hh>
#include <stdlib.h>
constexpr size_t buffer_size = (1024*1024);
constexpr int test_iteration_hstvismem = 5;
constexpr int test_iteration_noncohmem = 10;
constexpr int block_size = 512;
/**
* @addtogroup hipEventCreateWithFlags hipEventCreateWithFlags
* @{
* @ingroup EventTest
* `hipEventCreateWithFlags(hipEvent_t* event, unsigned flags)` -
* Create an event with the specified flags to control event behaviour.
* `hipEventCreateWithFlags (hipEvent_t *event, unsigned flags)` -
* begins graph capture on a stream
*/
/**
* Test Description
* ------------------------
* - Successfully create an event with all defined device flags.
* - Test simple event creation with hipEventCreateWithFlags api for each flag
* Test source
* ------------------------
* - unit/event/hipEventCreateWithFlags.cc
* - catch\unit\event\hipEventCreateWithFlags.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipEventCreateWithFlags_Positive") {
#if HT_AMD
const unsigned int flagUnderTest = GENERATE(hipEventDefault, hipEventBlockingSync, hipEventDisableTiming, hipEventInterprocess | hipEventDisableTiming, hipEventReleaseToDevice, hipEventReleaseToSystem);
const unsigned int flagUnderTest = GENERATE(hipEventDefault,
hipEventBlockingSync, hipEventDisableTiming,
hipEventInterprocess | hipEventDisableTiming,
hipEventReleaseToDevice, hipEventReleaseToSystem);
#else
// On Non-AMD platforms hipEventReleaseToDevice / hipEventReleaseToSystem are not defined
const unsigned int flagUnderTest = GENERATE(hipEventDefault, hipEventBlockingSync, hipEventDisableTiming, hipEventInterprocess | hipEventDisableTiming);
// On Non-AMD platforms hipEventReleaseToDevice / hipEventReleaseToSystem
// are not defined.
const unsigned int flagUnderTest = GENERATE(hipEventDefault,
hipEventBlockingSync,
hipEventDisableTiming,
hipEventInterprocess | hipEventDisableTiming);
#endif
hipEvent_t event;
HIP_CHECK(hipEventCreateWithFlags(&event, flagUnderTest));
REQUIRE(event != nullptr);
HIP_CHECK(hipEventDestroy(event));
}
/**
Since flags hipEventReleaseToSystem, hipEventDisableSystemFence and hipEventReleaseToDevice
are AMD specific flags, hence the following tests enabled only for AMD.
*/
#if HT_AMD
enum class eSyncToTest {
eStreamSynchronize,
eDeviceSynchronize,
eStreamWaitEvent,
eEventSynchronize
};
enum class eMemoryToTest {
eHostVisibleMemory,
eNonCoherentHostMemory,
eCoherentHostMemory
};
static void init_input(int* a, size_t size) {
unsigned int seed = time(nullptr);
for (size_t i = 0; i < size; i++) {
a[i] = (HipTest::RAND_R(&seed) & 0xFF);
}
}
static void check_output(int* inp, int* out, size_t size) {
for (size_t i = 0; i < size; i++) {
REQUIRE(out[i] == (inp[i]*inp[i]));
}
}
// local function
static void testMemCoherency(eSyncToTest test, eMemoryToTest mem,
uint32_t flags) {
hipDeviceProp_t prop;
HIP_CHECK(hipGetDeviceProperties(&prop, 0));
// If the GPU is not large bar then exit the test
if (prop.isLargeBar != 1) {
HipTest::HIP_SKIP_TEST("Skipping test as large bar is not supported");
return;
}
constexpr auto blocksPerCU = 6;
unsigned grid_size = HipTest::setNumBlocks(blocksPerCU,
block_size, buffer_size);
hipEvent_t event;
HIP_CHECK(hipEventCreateWithFlags(&event, flags));
hipStream_t stream;
HIP_CHECK(hipStreamCreateWithFlags(&stream, 0x0));
int *ibuf_h, *buf_d;
ibuf_h = new int[buffer_size];
REQUIRE(ibuf_h != nullptr);
int total_iter = 0;
if (mem == eMemoryToTest::eHostVisibleMemory) {
HIP_CHECK(hipMalloc(&buf_d, buffer_size*sizeof(int)));
total_iter = test_iteration_hstvismem;
} else if (mem == eMemoryToTest::eNonCoherentHostMemory) {
HIP_CHECK(hipHostMalloc(&buf_d, buffer_size*sizeof(int),
hipHostMallocNonCoherent));
total_iter = test_iteration_noncohmem;
} else if (mem == eMemoryToTest::eCoherentHostMemory) {
HIP_CHECK(hipHostMalloc(&buf_d, buffer_size*sizeof(int),
hipHostMallocCoherent));
total_iter = test_iteration_noncohmem;
}
for (int iter = 0; iter < total_iter; iter++) {
// Inititalize the buffer with random data
init_input(ibuf_h, buffer_size);
HIP_CHECK(hipMemcpy(buf_d, ibuf_h, sizeof(int)*buffer_size,
hipMemcpyDefault));
HipTest::vector_square<int><<<grid_size, block_size, 0, stream>>>(
buf_d, buf_d, buffer_size);
HIP_CHECK(hipEventRecord(event, stream));
// test different synchronization APIs
if (test == eSyncToTest::eStreamSynchronize) {
HIP_CHECK(hipStreamSynchronize(stream));
} else if (test == eSyncToTest::eDeviceSynchronize) {
HIP_CHECK(hipDeviceSynchronize());
} else if (test == eSyncToTest::eEventSynchronize) {
HIP_CHECK(hipEventSynchronize(event));
} else if (test == eSyncToTest::eStreamWaitEvent) {
HIP_CHECK(hipStreamWaitEvent(stream, event, 0));
}
check_output(ibuf_h, buf_d, buffer_size);
}
delete[] ibuf_h;
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipEventDestroy(event));
if (mem == eMemoryToTest::eHostVisibleMemory) {
HIP_CHECK(hipFree(buf_d));
} else if ((mem == eMemoryToTest::eNonCoherentHostMemory) ||
(mem == eMemoryToTest::eCoherentHostMemory)) {
HIP_CHECK(hipHostFree(buf_d));
}
}
/**
* Test Description
* ------------------------
* - Check Synchronization effect on Host Visible Memory.
* Disable System fence when creating an event. Create a chunk of Host Visisble
* Memory using hipMalloc and initialize the memory with user data. Launch a
* kernel that writes to this memory location. Perform different synchronizations
* and validate that updated values are seen from host.
* Test source
* ------------------------
* - catch\unit\event\hipEventCreateWithFlags.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_hipEventCreateWithFlags_DisableSystemFence_HstVisMem") {
SECTION("Check with hipStreamSynchronize") {
INFO("Check with hipStreamSynchronize");
testMemCoherency(eSyncToTest::eStreamSynchronize,
eMemoryToTest::eHostVisibleMemory, hipEventDisableSystemFence);
}
SECTION("Check with hipDeviceSynchronize") {
INFO("Check with hipDeviceSynchronize");
testMemCoherency(eSyncToTest::eDeviceSynchronize,
eMemoryToTest::eHostVisibleMemory, hipEventDisableSystemFence);
}
}
/**
* Test Description
* ------------------------
* - Check Synchronization effect on Host Visible Memory.
* Use Default Flag when creating an event. Create a chunk of Host Visisble
* Memory using hipMalloc and initialize the memory with user data. Launch
* a kernel that writes to this memory location. Perform different synchronizations
* and validate that updated values are seen from host.
* ------------------------
* - catch\unit\event\hipEventCreateWithFlags.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_hipEventCreateWithFlags_DefaultFlg_HstVisMem") {
SECTION("Check with hipStreamSynchronize") {
INFO("Check with hipStreamSynchronize");
testMemCoherency(eSyncToTest::eStreamSynchronize,
eMemoryToTest::eHostVisibleMemory, hipEventDefault);
}
SECTION("Check with hipDeviceSynchronize") {
INFO("Check with hipDeviceSynchronize");
testMemCoherency(eSyncToTest::eDeviceSynchronize,
eMemoryToTest::eHostVisibleMemory, hipEventDefault);
}
SECTION("Check with hipEventSynchronize") {
INFO("Check with hipEventSynchronize");
testMemCoherency(eSyncToTest::eEventSynchronize,
eMemoryToTest::eHostVisibleMemory, hipEventDefault);
}
}
/**
* Test Description
* ------------------------
* - Check Synchronization effect on Non Coherent Host Memory.
* Disable System fence when creating an event. Create a chunk of Host Visisble
* Memory using hipMalloc and initialize the memory with user data. Launch a
* kernel that writes to this memory location. Perform different synchronizations
* and validate that updated values are seen from host.
* ------------------------
* - catch\unit\event\hipEventCreateWithFlags.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_hipEventCreateWithFlags_DisableSystemFence_NonCohHstMem") {
SECTION("Check with hipStreamSynchronize") {
INFO("Check with hipStreamSynchronize");
testMemCoherency(eSyncToTest::eStreamSynchronize,
eMemoryToTest::eNonCoherentHostMemory, hipEventDisableSystemFence);
}
SECTION("Check with hipDeviceSynchronize") {
INFO("Check with hipDeviceSynchronize");
testMemCoherency(eSyncToTest::eDeviceSynchronize,
eMemoryToTest::eNonCoherentHostMemory, hipEventDisableSystemFence);
}
}
/**
* Test Description
* ------------------------
* - Check Synchronization effect on Non Coherent Host Memory.
* Use Default Flag when creating an event. Create a chunk of Host Visisble
* Memory using hipMalloc and initialize the memory with user data. Launch
* a kernel that writes to this memory location. Perform different
* synchronizations and validate that updated values are seen from host.
* ------------------------
* - catch\unit\event\hipEventCreateWithFlags.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_hipEventCreateWithFlags_DefaultFlg_NonCohHstMem") {
SECTION("Check with hipStreamSynchronize") {
INFO("Check with hipStreamSynchronize");
testMemCoherency(eSyncToTest::eStreamSynchronize,
eMemoryToTest::eNonCoherentHostMemory, hipEventDefault);
}
SECTION("Check with hipDeviceSynchronize") {
INFO("Check with hipDeviceSynchronize");
testMemCoherency(eSyncToTest::eDeviceSynchronize,
eMemoryToTest::eNonCoherentHostMemory, hipEventDefault);
}
SECTION("Check with hipEventSynchronize") {
INFO("Check with hipEventSynchronize");
testMemCoherency(eSyncToTest::eEventSynchronize,
eMemoryToTest::eNonCoherentHostMemory, hipEventDefault);
}
}
/**
* Test Description
* ------------------------
* - Check Synchronization effect on Coherent Host Memory.
* Disable System fence when creating an event. Create a chunk of Host Visisble
* Memory using hipMalloc and initialize the memory with user data. Launch a
* kernel that writes to this memory location. Perform different synchronizations
* and validate that updated values are seen from host.
* ------------------------
* - catch\unit\event\hipEventCreateWithFlags.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_hipEventCreateWithFlags_DisableSystemFence_CohHstMem") {
SECTION("Check with hipStreamSynchronize") {
INFO("Check with hipStreamSynchronize");
testMemCoherency(eSyncToTest::eStreamSynchronize,
eMemoryToTest::eCoherentHostMemory, hipEventDisableSystemFence);
}
SECTION("Check with hipDeviceSynchronize") {
INFO("Check with hipDeviceSynchronize");
testMemCoherency(eSyncToTest::eDeviceSynchronize,
eMemoryToTest::eCoherentHostMemory, hipEventDisableSystemFence);
}
SECTION("Check with hipEventSynchronize") {
INFO("Check with hipEventSynchronize");
testMemCoherency(eSyncToTest::eEventSynchronize,
eMemoryToTest::eCoherentHostMemory, hipEventDisableSystemFence);
}
}
/**
* Test Description
* ------------------------
* - Check Synchronization effect on Coherent Host Memory.
* Use Default Flag when creating an event. Create a chunk of Host Visisble
* Memory using hipMalloc and initialize the memory with user data. Launch a
* kernel that writes to this memory location. Perform different synchronizations
* and validate that updated values are seen from host.
* ------------------------
* - catch\unit\event\hipEventCreateWithFlags.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_hipEventCreateWithFlags_DefaultFlg_CohHstMem") {
SECTION("Check with hipStreamSynchronize") {
INFO("Check with hipStreamSynchronize");
testMemCoherency(eSyncToTest::eStreamSynchronize,
eMemoryToTest::eCoherentHostMemory, hipEventDefault);
}
SECTION("Check with hipDeviceSynchronize") {
INFO("Check with hipDeviceSynchronize");
testMemCoherency(eSyncToTest::eDeviceSynchronize,
eMemoryToTest::eCoherentHostMemory, hipEventDefault);
}
SECTION("Check with hipEventSynchronize") {
INFO("Check with hipEventSynchronize");
testMemCoherency(eSyncToTest::eEventSynchronize,
eMemoryToTest::eCoherentHostMemory, hipEventDefault);
}
}
#endif