SWDEV-289401 - Catch2 tests for hip event APIs

Change-Id: I7c28f842282e07c21656fb92ddbb1a9ad32d752c


[ROCm/hip commit: 5816d20752]

projects/hip/tests/catch/hipTestMain/CMakeLists.txt

@@ -12,6 +12,7 @@ endif()
 target_link_libraries(UnitTests PRIVATE UnitDeviceTests
                                         MemoryTest
                                         StreamTest
+                                        EventTest
                                         OccupancyTest
                                         DeviceTest
                                         RTC

projects/hip/tests/catch/include/hip_test_common.hh

@@ -49,6 +49,11 @@ THE SOFTWARE.
 #define HIP_ASSERT(x)                                                                              \
   { REQUIRE((x)); }
 
+#ifdef __cplusplus
+  #include <iostream>
+  #include <iomanip>
+  #include <chrono>
+#endif
 
 // Utility Functions
 namespace HipTest {
@@ -57,4 +62,28 @@ static inline int getDeviceCount() {
   HIP_CHECK(hipGetDeviceCount(&dev));
   return dev;
 }
+
+// Returns the current system time in microseconds
+static inline long long get_time() {
+  return std::chrono::high_resolution_clock::now().time_since_epoch()
+      /std::chrono::microseconds(1);
+}
+
+static inline double elapsed_time(long long startTimeUs, long long stopTimeUs) {
+  return ((double)(stopTimeUs - startTimeUs)) / ((double)(1000));
+}
+
+static inline unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N) {
+  int device;
+  HIP_CHECK(hipGetDevice(&device));
+  hipDeviceProp_t props;
+  HIP_CHECK(hipGetDeviceProperties(&props, device));
+
+  unsigned blocks = props.multiProcessorCount * blocksPerCU;
+  if (blocks * threadsPerBlock > N) {
+    blocks = (N + threadsPerBlock - 1) / threadsPerBlock;
+  }
+
+  return blocks;
+}
 }

projects/hip/tests/catch/multiproc/hipMallocConcurrency.cc

@@ -10,29 +10,12 @@
 #include <limits>
 #include <atomic>
 
-
 #include <hip_test_common.hh>
 
 size_t N = 4 * 1024 * 1024;
 unsigned blocksPerCU = 6;  // to hide latency
 unsigned threadsPerBlock = 256;
 
-
-unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N) {
-  int device;
-  HIP_CHECK(hipGetDevice(&device));
-  hipDeviceProp_t props;
-  HIP_CHECK(hipGetDeviceProperties(&props, device));
-
-  unsigned blocks = props.multiProcessorCount * blocksPerCU;
-  if (blocks * threadsPerBlock > N) {
-    blocks = (N + threadsPerBlock - 1) / threadsPerBlock;
-  }
-
-  return blocks;
-}
-
-
 /**
  * Validates data consitency on supplied gpu
  */
@@ -48,7 +31,7 @@ bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
   printf("tgs allocating..\n");
   HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
 
-  unsigned blocks = setNumBlocks(blocksPerCU, threadsPerBlock, N);
+  unsigned blocks =  HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
 
   HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
   HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));

projects/hip/tests/catch/unit/CMakeLists.txt

@@ -1,6 +1,7 @@
 add_subdirectory(memory)
 add_subdirectory(deviceLib)
 add_subdirectory(stream)
+add_subdirectory(event)
 add_subdirectory(occupancy)
 add_subdirectory(device)
 add_subdirectory(rtc)

projects/hip/tests/catch/unit/event/CMakeLists.txt

@@ -0,0 +1,14 @@
+# Common Tests - Test independent of all platforms
+set(TEST_SRC
+    Unit_hipEvent_Negative.cc
+    Unit_hipEvent.cc
+    Unit_hipEventElapsedTime.cc
+    Unit_hipEventRecord.cc
+    Unit_hipEventIpc.cc
+)
+
+# Create shared lib of all tests
+add_library(EventTest SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
+
+# Add dependency on build_tests to build it on this custom target
+add_dependencies(build_tests EventTest)

projects/hip/tests/catch/unit/event/Unit_hipEvent.cc

@@ -0,0 +1,187 @@
+/*
+Copyright (c) 2021 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
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+// Tests hipEventRecord and hipEventElapsedTime with different scenarios
+// and confirms if these are working as expected
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+#include <hip_test_common.hh>
+
+int tests = -1;
+enum SyncMode {
+    syncNone,
+    syncStream,
+    syncStopEvent,
+};
+
+const char* syncModeString(int syncMode) {
+    switch (syncMode) {
+        case syncNone:
+            return "syncNone";
+        case syncStream:
+            return "syncStream";
+        case syncStopEvent:
+            return "syncStopEvent";
+        default:
+            return "unknown";
+    };
+};
+
+void test(unsigned testMask, int* C_d, int* C_h, int64_t numElements, hipStream_t stream,
+          int waitStart, SyncMode syncMode) {
+    if (!(testMask & tests)) {
+        return;
+    }
+    std::cout << "\n test 0x " << testMask << ": stream= " << stream << " waitStart= " << waitStart
+              << " syncMode= " << syncModeString(syncMode) << std::endl;
+
+    size_t sizeBytes = numElements * sizeof(int);
+
+    int count = 100;
+    int init0 = 0;
+    HIP_CHECK(hipMemset(C_d, init0, sizeBytes));
+    for (int i = 0; i < numElements; i++) {
+        C_h[i] = -1;  // initialize
+    }
+
+    hipEvent_t neverCreated = 0, neverRecorded, timingDisabled;
+    HIP_CHECK(hipEventCreate(&neverRecorded));
+    HIP_CHECK(hipEventCreateWithFlags(&timingDisabled, hipEventDisableTiming));
+
+    hipEvent_t start, stop;
+    HIP_CHECK(hipEventCreate(&start));
+    HIP_CHECK(hipEventCreate(&stop));
+    unsigned blocksPerCU = 6;
+    unsigned threadsPerBlock = 256;
+
+    unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
+
+    HIP_CHECK(hipEventRecord(timingDisabled, stream));
+    // sandwhich a kernel:
+    HIP_CHECK(hipEventRecord(start, stream));
+    hipLaunchKernelGGL(HipTest::addCountReverse, dim3(blocks), dim3(threadsPerBlock), 0, stream,
+                       static_cast<const int*>(C_d), C_h, numElements, count);
+    HIP_CHECK(hipEventRecord(stop, stream));
+
+    if (waitStart) {
+        HIP_CHECK(hipEventSynchronize(start));
+    }
+
+    hipError_t expectedStopError = hipSuccess;
+
+    // How to wait for the events to finish:
+    switch (syncMode) {
+        case syncNone:
+            expectedStopError = hipErrorNotReady;
+            break;
+        case syncStream:
+            HIP_CHECK(hipStreamSynchronize(stream));  // wait for recording to finish...
+            break;
+        case syncStopEvent:
+            HIP_CHECK(hipEventSynchronize(stop));
+            break;
+        default:
+            assert(0);
+    };
+
+    float t;
+
+    hipError_t e = hipEventElapsedTime(&t, start, start);
+    if ((e != hipSuccess) && (e != hipErrorNotReady || syncMode != syncNone)) {
+        printf("start event not in expected state, was %d=%s\n", e, hipGetErrorName(e));
+        REQUIRE(false);
+    }
+
+    if (e == hipSuccess) assert(t == 0.0f);
+
+    // stop usually ready unless we skipped the synchronization (syncNone)
+    HIP_ASSERT(hipEventElapsedTime(&t, stop, stop) == expectedStopError);
+    if (e == hipSuccess) assert(t == 0.0f);
+
+    e = hipEventElapsedTime(&t, start, stop);
+    HIP_ASSERT(e == expectedStopError);
+    if (expectedStopError == hipSuccess) assert(t > 0.0f);
+    printf("time=%6.2f error=%s\n", t, hipGetErrorName(e));
+
+    e = hipEventElapsedTime(&t, stop, start);
+    HIP_ASSERT(e == expectedStopError);
+    if (expectedStopError == hipSuccess) assert(t < 0.0f);
+    printf("negtime=%6.2f error=%s\n", t, hipGetErrorName(e));
+
+    {
+        // Check some error conditions for incomplete events:
+        HIP_ASSERT(hipEventElapsedTime(&t, timingDisabled, stop) == hipErrorInvalidHandle);
+        HIP_ASSERT(hipEventElapsedTime(&t, start, timingDisabled) == hipErrorInvalidHandle);
+
+        HIP_ASSERT(hipEventElapsedTime(&t, neverCreated, stop) == hipErrorInvalidHandle);
+        HIP_ASSERT(hipEventElapsedTime(&t, start, neverCreated) == hipErrorInvalidHandle);
+
+        HIP_ASSERT(hipEventElapsedTime(&t, neverRecorded, stop) == hipErrorInvalidHandle);
+        HIP_ASSERT(hipEventElapsedTime(&t, start, neverRecorded) == hipErrorInvalidHandle);
+    }
+
+    HIP_CHECK(hipEventDestroy(neverRecorded));
+    HIP_CHECK(hipEventDestroy(timingDisabled));
+
+    HIP_CHECK(hipEventDestroy(start));
+    HIP_CHECK(hipEventDestroy(stop));
+
+    // Clear out everything:
+    HIP_CHECK(hipDeviceSynchronize());
+
+    printf("test:   OK  \n");
+}
+
+void runTests(int64_t numElements) {
+    size_t sizeBytes = numElements * sizeof(int);
+
+    printf("test: starting sequence with sizeBytes=%zu bytes, %6.2f MB\n", sizeBytes,
+           sizeBytes / 1024.0 / 1024.0);
+
+
+    int *C_h, *C_d;
+    HIP_CHECK(hipMalloc(&C_d, sizeBytes));
+    HIP_CHECK(hipHostMalloc(&C_h, sizeBytes));
+
+    hipStream_t stream;
+    HIP_CHECK(hipStreamCreateWithFlags(&stream, 0x0));
+
+    for (int waitStart = 1; waitStart >= 0; waitStart--) {
+        unsigned W = waitStart ? 0x1000 : 0;
+        test(W | 0x01, C_d, C_h, numElements, 0, 0, syncNone);
+        test(W | 0x02, C_d, C_h, numElements, stream, 0, syncNone);
+        test(W | 0x04, C_d, C_h, numElements, 0, waitStart, syncStream);
+        test(W | 0x08, C_d, C_h, numElements, stream, waitStart, syncStream);
+        test(W | 0x10, C_d, C_h, numElements, 0, waitStart, syncStopEvent);
+        test(W | 0x20, C_d, C_h, numElements, stream, waitStart, syncStopEvent);
+    }
+
+
+    HIP_CHECK(hipStreamDestroy(stream));
+    HIP_CHECK(hipFree(C_d));
+    HIP_CHECK(hipHostFree(C_h));
+}
+
+
+TEST_CASE("Unit_hipEvent") {
+  runTests(80000000);
+}

projects/hip/tests/catch/unit/event/Unit_hipEventElapsedTime.cc

@@ -0,0 +1,100 @@
+/*
+Copyright (c) 2021 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
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <iostream>
+
+TEST_CASE("Unit_hipEventElapsedTime_NullCheck") {
+  hipEvent_t start = nullptr, end = nullptr;
+  float tms = 1.0f;
+  HIP_ASSERT(hipEventElapsedTime(nullptr, start, end) == hipErrorInvalidValue);
+#ifndef __HIP_PLATFORM_NVIDIA__
+  // On NVCC platform API throws seg fault hence skipping
+  HIP_ASSERT(hipEventElapsedTime(&tms, nullptr, end) == hipErrorInvalidHandle);
+  HIP_ASSERT(hipEventElapsedTime(&tms, start, nullptr) == hipErrorInvalidHandle);
+#endif
+}
+
+TEST_CASE("Unit_hipEventElapsedTime_DisableTiming") {
+  float timeElapsed = 1.0f;
+  hipEvent_t start, stop;
+  HIP_CHECK(hipEventCreateWithFlags(&start, hipEventDisableTiming));
+  HIP_CHECK(hipEventCreateWithFlags(&stop, hipEventDisableTiming));
+  HIP_ASSERT(hipEventElapsedTime(&timeElapsed, start, stop) == hipErrorInvalidHandle);
+  HIP_CHECK(hipEventDestroy(start));
+  HIP_CHECK(hipEventDestroy(stop));
+}
+
+TEST_CASE("Unit_hipEventElapsedTime_DifferentDevices") {
+  int devCount = 0;
+  HIP_CHECK(hipGetDeviceCount(&devCount));
+  if (devCount > 1) {
+    // create event on dev=0
+    HIP_CHECK(hipSetDevice(0));
+    hipEvent_t start;
+    hipEvent_t start1;
+    HIP_CHECK(hipEventCreate(&start));
+    HIP_CHECK(hipEventCreate(&start1));
+
+    HIP_CHECK(hipEventRecord(start, nullptr));
+    HIP_CHECK(hipEventSynchronize(start));
+
+    // create event on dev=1
+    HIP_CHECK(hipSetDevice(1));
+    hipEvent_t stop;
+    HIP_CHECK(hipEventCreate(&stop));
+
+    // start1 on device 0 but null stream on device 1
+    HIP_ASSERT(hipEventRecord(start1, nullptr) == hipErrorInvalidHandle);
+
+    HIP_CHECK(hipEventRecord(stop, nullptr));
+    HIP_CHECK(hipEventSynchronize(stop));
+
+    float tElapsed = 1.0f;
+    // start on device 0 but stop on device 1
+    HIP_ASSERT(hipEventElapsedTime(&tElapsed,start,stop) == hipErrorInvalidHandle);
+
+    HIP_CHECK(hipEventDestroy(start));
+    HIP_CHECK(hipEventDestroy(start1));
+    HIP_CHECK(hipEventDestroy(stop));
+  }
+}
+
+TEST_CASE("Unit_hipEventElapsedTime") {
+  hipEvent_t start;
+  HIP_CHECK(hipEventCreate(&start));
+
+  hipEvent_t stop;
+  HIP_CHECK(hipEventCreate(&stop));
+
+  HIP_CHECK(hipEventRecord(start, nullptr));
+  HIP_CHECK(hipEventSynchronize(start));
+
+  HIP_CHECK(hipEventRecord(stop, nullptr));
+  HIP_CHECK(hipEventSynchronize(stop));
+
+  float tElapsed = 1.0f;
+  HIP_CHECK(hipEventElapsedTime(&tElapsed, start, stop));
+
+  HIP_CHECK(hipEventDestroy(start));
+  HIP_CHECK(hipEventDestroy(stop));
+}

projects/hip/tests/catch/unit/event/Unit_hipEventIpc.cc

@@ -0,0 +1,104 @@
+/*
+Copyright (c) 2021 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
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+// Test hipEventIpc behavior.
+
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+#include <hip_test_common.hh>
+
+
+TEST_CASE("Unit_hipEventIpc") {
+    size_t N = 4 * 1024 * 1024;
+    unsigned threadsPerBlock = 256;
+    int iterations = 1;
+
+    unsigned blocks = (N + threadsPerBlock - 1) / threadsPerBlock;
+    if (blocks > 1024) blocks = 1024;
+    if (blocks == 0) blocks = 1;
+
+    printf("N=%zu (A+B+C= %6.1f MB total) blocks=%u threadsPerBlock=%u iterations=%d\n", N,
+           ((double)3 * N * sizeof(float)) / 1024 / 1024, blocks, threadsPerBlock, iterations);
+    printf("iterations=%d\n", iterations);
+
+    size_t Nbytes = N * sizeof(float);
+
+    float *A_h, *B_h, *C_h;
+    float *A_d, *B_d, *C_d;
+    HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N);
+
+    hipEvent_t start, stop;
+
+    // NULL stream check:
+    HIP_CHECK(hipEventCreateWithFlags(&start, hipEventDisableTiming|hipEventInterprocess));
+    HIP_CHECK(hipEventCreateWithFlags(&stop, hipEventDisableTiming|hipEventInterprocess));
+
+    HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+    HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
+
+
+    for (int i = 0; i < iterations; i++) {
+        //--- START TIMED REGION
+        long long hostStart = HipTest::get_time();
+        // Record the start event
+        HIP_CHECK(hipEventRecord(start, NULL));
+
+        hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
+                        static_cast<const float*>(A_d), static_cast<const float*>(B_d), C_d, N);
+
+
+        HIP_CHECK(hipEventRecord(stop, NULL));
+        HIP_CHECK(hipEventSynchronize(stop));
+        HIP_CHECK(hipEventQuery(stop));
+        long long hostStop = HipTest::get_time();
+        //--- STOP TIMED REGION
+
+
+        float eventMs = 1.0f;
+        // should fail due to hipEventDisableTiming
+        REQUIRE(hipSuccess != hipEventElapsedTime(&eventMs, start, stop));
+        float hostMs = HipTest::elapsed_time(hostStart, hostStop);
+
+        printf("host_time (chrono)                =%6.3fms\n", hostMs);
+        printf("kernel_time (hipEventElapsedTime) =%6.3fms\n", eventMs);
+        printf("\n");
+
+    }
+
+    hipIpcEventHandle_t ipc_handle;
+    HIP_CHECK(hipIpcGetEventHandle(&ipc_handle, start));
+
+    hipEvent_t ipc_event;
+    hipError_t err = hipIpcOpenEventHandle(&ipc_event, ipc_handle);
+
+    // hipIpcOpenEventHandle() should be called in a different process, hence it should fail here
+    REQUIRE(err == hipErrorInvalidContext);
+
+    HIP_CHECK(hipEventDestroy(start));
+    HIP_CHECK(hipEventDestroy(stop));
+
+    HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
+
+    HipTest::checkVectorADD(A_h, B_h, C_h, N, true);
+
+}

projects/hip/tests/catch/unit/event/Unit_hipEventRecord.cc

@@ -0,0 +1,91 @@
+/*
+Copyright (c) 2021 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
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+// Test hipEventRecord serialization behavior.
+// Through manual inspection of the reported timestamps, can determine if recording a NULL event
+// forces synchronization : set
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+#include <hip_test_common.hh>
+
+TEST_CASE("Unit_hipEventRecord") {
+    size_t N = 4 * 1024 * 1024;
+    unsigned threadsPerBlock = 256;
+    int iterations = 1;
+
+    unsigned blocks = (N + threadsPerBlock - 1) / threadsPerBlock;
+    if (blocks > 1024) blocks = 1024;
+    if (blocks == 0) blocks = 1;
+
+    printf("N=%zu (A+B+C= %6.1f MB total) blocks=%u threadsPerBlock=%u iterations=%d\n", N,
+           ((double)3 * N * sizeof(float)) / 1024 / 1024, blocks, threadsPerBlock, iterations);
+    printf("iterations=%d\n", iterations);
+
+    size_t Nbytes = N * sizeof(float);
+
+    float *A_h, *B_h, *C_h;
+    float *A_d, *B_d, *C_d;
+    HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N);
+
+    hipEvent_t start, stop;
+
+    // NULL stream check:
+    HIP_CHECK(hipEventCreate(&start));
+    HIP_CHECK(hipEventCreate(&stop));
+
+    HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+    HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
+
+    for (int i = 0; i < iterations; i++) {
+        //--- START TIMED REGION
+        long long hostStart = HipTest::get_time();
+        // Record the start event
+        HIP_CHECK(hipEventRecord(start, NULL));
+
+        hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
+                        static_cast<const float*>(A_d), static_cast<const float*>(B_d), C_d, N);
+
+        HIP_CHECK(hipEventRecord(stop, NULL));
+        HIP_CHECK(hipEventSynchronize(stop));
+        long long hostStop = HipTest::get_time();
+        //--- STOP TIMED REGION
+
+        float eventMs = 1.0f;
+        HIP_CHECK(hipEventElapsedTime(&eventMs, start, stop));
+        float hostMs = HipTest::elapsed_time(hostStart, hostStop);
+
+        printf("host_time (chrono)                =%6.3fms\n", hostMs);
+        printf("kernel_time (hipEventElapsedTime) =%6.3fms\n", eventMs);
+        printf("\n");
+
+        // Make sure timer is timing something...
+        REQUIRE(eventMs > 0.0f);
+    }
+
+    HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
+
+    HIP_CHECK(hipEventDestroy(start));
+    HIP_CHECK(hipEventDestroy(stop));
+
+    HipTest::checkVectorADD(A_h, B_h, C_h, N, true);
+
+}

projects/hip/tests/catch/unit/event/Unit_hipEvent_Negative.cc

@@ -0,0 +1,48 @@
+/*
+Copyright (c) 2021 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
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+
+TEST_CASE("Unit_hipEventCreate_NullCheck") {
+  hipEvent_t start_event;
+  auto res = hipEventCreate(nullptr);
+  REQUIRE(res != hipSuccess);
+  res = hipEventCreateWithFlags(nullptr, 0);
+  REQUIRE(res != hipSuccess);
+  res = hipEventCreateWithFlags(&start_event, 10);
+  REQUIRE(res != hipSuccess);
+}
+
+TEST_CASE("Unit_hipEventSynchronize_NullCheck") {
+  auto res = hipEventSynchronize(nullptr);
+  REQUIRE(res != hipSuccess);
+}
+
+TEST_CASE("Unit_hipEventQuery_NullCheck") {
+  auto res = hipEventQuery(nullptr);
+  REQUIRE(res != hipSuccess);
+}
+
+TEST_CASE("Unit_hipEventDestroy_NullCheck") {
+  auto res = hipEventDestroy(nullptr);
+  REQUIRE(res != hipSuccess);
+}