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

Change-Id: Ied3dce437f34d793eb3715999125924ff74ff9b8

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",

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",

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
  * @{

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 <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)` -
+ * `hipEventCreateWithFlags (hipEvent_t *event, unsigned flags)` -
- * Create an event with the specified flags to control event behaviour.
+ * 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
+  // On Non-AMD platforms hipEventReleaseToDevice / hipEventReleaseToSystem
-  const unsigned int flagUnderTest = GENERATE(hipEventDefault, hipEventBlockingSync, hipEventDisableTiming, hipEventInterprocess | hipEventDisableTiming);
+  // 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