From 23b02df2a188f02fd104f4d585de5b1de0d350c8 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 14 Aug 2023 20:53:37 +0530 Subject: [PATCH] SWDEV-385489 - [catch2][dtest] Adding scenarios for hipEventCreateWithFlags() with hipEventDisableSystemFence and default flags (#282) Change-Id: Ied3dce437f34d793eb3715999125924ff74ff9b8 --- .../config/config_amd_windows_MI2xx.json | 7 + .../config/config_amd_windows_common.json | 6 + catch/include/hip_test_defgroups.hh | 7 + catch/unit/event/hipEventCreateWithFlags.cc | 314 +++++++++++++++++- 4 files changed, 323 insertions(+), 11 deletions(-) diff --git a/catch/hipTestMain/config/config_amd_windows_MI2xx.json b/catch/hipTestMain/config/config_amd_windows_MI2xx.json index 92ce9593a5..27ab1dd12b 100644 --- a/catch/hipTestMain/config/config_amd_windows_MI2xx.json +++ b/catch/hipTestMain/config/config_amd_windows_MI2xx.json @@ -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", diff --git a/catch/hipTestMain/config/config_amd_windows_common.json b/catch/hipTestMain/config/config_amd_windows_common.json index 405556ecc1..4f76812ef6 100644 --- a/catch/hipTestMain/config/config_amd_windows_common.json +++ b/catch/hipTestMain/config/config_amd_windows_common.json @@ -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", diff --git a/catch/include/hip_test_defgroups.hh b/catch/include/hip_test_defgroups.hh index 119b095dcc..7d6fd3f731 100644 --- a/catch/include/hip_test_defgroups.hh +++ b/catch/include/hip_test_defgroups.hh @@ -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 * @{ diff --git a/catch/unit/event/hipEventCreateWithFlags.cc b/catch/unit/event/hipEventCreateWithFlags.cc index c758e12a55..875d7f4295 100644 --- a/catch/unit/event/hipEventCreateWithFlags.cc +++ b/catch/unit/event/hipEventCreateWithFlags.cc @@ -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 +#include +#include +#include + +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<<>>( + 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