From 449cc5096efe1520e2625cac14f150fb3dbf8f79 Mon Sep 17 00:00:00 2001
From: kjayapra-amd <karthik.jayaprakash@amd.com>
Date: Mon, 8 Mar 2021 20:44:30 -0500
Subject: [PATCH] SWDEV-276169 - Multiprocess IPC tests for Events and Memory.

Change-Id: I4a7af60e732de234a535574aa8597a7abd0b899b


[ROCm/clr commit: 57615530a7e496549e85f36e78e3611ee865c692]
---
 projects/clr/hipamd/bin/hipcc                 |   2 +-
 .../clr/hipamd/tests/src/ipc/MultiProcess.h   | 157 ++++++++++++++++++
 .../tests/src/ipc/hipMultiProcIpcEvent.cpp    | 126 ++++++++++++++
 .../tests/src/ipc/hipMultiProcIpcMem.cpp      |  98 +++++++++++
 .../{hipSimpleIpc.cpp => hipSimpleIpcMem.cpp} |   0
 projects/clr/hipamd/tests/src/test_common.cpp |   5 +
 projects/clr/hipamd/tests/src/test_common.h   |   1 +
 7 files changed, 388 insertions(+), 1 deletion(-)
 create mode 100755 projects/clr/hipamd/tests/src/ipc/MultiProcess.h
 create mode 100755 projects/clr/hipamd/tests/src/ipc/hipMultiProcIpcEvent.cpp
 create mode 100755 projects/clr/hipamd/tests/src/ipc/hipMultiProcIpcMem.cpp
 rename projects/clr/hipamd/tests/src/ipc/{hipSimpleIpc.cpp => hipSimpleIpcMem.cpp} (100%)

diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc
index 7509f14988..17a80755db 100755
--- a/projects/clr/hipamd/bin/hipcc
+++ b/projects/clr/hipamd/bin/hipcc
@@ -688,7 +688,7 @@ if ($HIP_PLATFORM eq "amd") {
         $HIPCXXFLAGS .= " -fhip-new-launch-api";
     }
     if (not $isWindows) {
-        $HIPLDFLAGS .= " -lgcc_s -lgcc -lpthread -lm";
+        $HIPLDFLAGS .= " -lgcc_s -lgcc -lpthread -lm -lrt";
     }
 
     if (not $isWindows  and not $compileOnly) {
diff --git a/projects/clr/hipamd/tests/src/ipc/MultiProcess.h b/projects/clr/hipamd/tests/src/ipc/MultiProcess.h
new file mode 100755
index 0000000000..f996887ed5
--- /dev/null
+++ b/projects/clr/hipamd/tests/src/ipc/MultiProcess.h
@@ -0,0 +1,157 @@
+#pragma once
+
+#ifdef __unix__
+
+#include <string>
+#include <atomic>
+#include <unistd.h>
+#include <fcntl.h>
+#include <sys/mman.h>
+
+template <typename T>
+struct Shmem {
+  std::atomic<T> handle_;
+  std::atomic<int> done_counter_;
+};
+
+template <typename T>
+struct ShmemMeta {
+  std::string shmem_name_;
+  int shmem_fd_;
+  Shmem<T>* shmem_;
+};
+
+template <typename T>
+class MultiProcess {
+public:
+  MultiProcess(size_t num_proc) : num_proc_(num_proc) {}
+  ~MultiProcess();
+
+  void DebugInfo(pid_t pid);
+
+  pid_t SpawnProcess(bool debug_bkpt);
+  bool CreateShmem();
+
+  bool WriteHandleToShmem(T ipc_handle);
+  bool WaitTillAllChildReads();
+
+  bool ReadHandleFromShmem(T& ipc_handle);
+  bool NotifyParentDone();
+
+private:
+  const size_t num_proc_;
+  bool debug_proc_;
+  ShmemMeta<T> shmem_meta_obj_;
+};
+
+// Template Implementations
+template <typename T>
+MultiProcess<T>::~MultiProcess() {
+  if(munmap(shmem_meta_obj_.shmem_, sizeof(Shmem<T>)) < 0) {
+    std::cout<<"Error Unmapping shared memory "<<std::endl;
+    exit(0);
+  }
+}
+
+template <typename T>
+void MultiProcess<T>::DebugInfo(pid_t pid) {
+  const int delay = 1;
+
+  if (pid == 0) {
+    std::cout<<" Child Process with ID: "<<getpid()<<std::endl;
+  } else {
+    std::cout<<" Parent Process with ID: "<<getpid()<<std::endl;
+  }
+
+  volatile int flag = 0;
+  while (!flag) {
+    sleep(delay);
+  }
+}
+
+template <typename T>
+pid_t MultiProcess<T>::SpawnProcess(bool debug_bkpt) {
+  if (num_proc_ < 0) {
+    std::cout<<"Num Process cannot be less than 1"<<std::endl;
+    return -1;
+  }
+
+  pid_t pid;
+  for (size_t proc_idx = 0; proc_idx < num_proc_; ++proc_idx) {
+    pid = fork();
+    if (pid < 0) {
+      std::cout<<"Fork Failed"<<std::endl;
+      assert(false);
+    } else if (pid == 0) {
+      //Child Process, so break
+      break;
+    }
+  }
+
+  if (debug_bkpt) {
+    DebugInfo(pid);
+  }
+
+  return pid;
+}
+
+template <typename T>
+bool MultiProcess<T>::CreateShmem() {
+  if (num_proc_ < 0) {
+    std::cout<<"Num Process cannot be less than 1"<<std::endl;
+    return false;
+  }
+
+  char name_template[] = "/tmp/eventXXXXX";
+  int temp_fd = mkstemp(name_template);
+  shmem_meta_obj_.shmem_name_ = name_template;
+  shmem_meta_obj_.shmem_name_.replace(0, 5, "/hip_");
+  shmem_meta_obj_.shmem_fd_ = shm_open(shmem_meta_obj_.shmem_name_.c_str(),
+                                       O_RDWR | O_CREAT, 0777);
+
+  if (ftruncate(shmem_meta_obj_.shmem_fd_, sizeof(ShmemMeta<T>)) != 0) {
+    std::cout<<"Cannot FTruncate "<<std::endl;
+    exit(0);
+  }
+
+  shmem_meta_obj_.shmem_ = (Shmem<T>*)mmap(0, sizeof(Shmem<T>), PROT_READ | PROT_WRITE,
+                                           MAP_SHARED, shmem_meta_obj_.shmem_fd_, 0);
+  memset(&shmem_meta_obj_.shmem_->handle_, 0x00, sizeof(T));
+  shmem_meta_obj_.shmem_->done_counter_ = -1;
+
+  return true;
+}
+
+template <typename T>
+bool MultiProcess<T>::WriteHandleToShmem(T ipc_handle) {
+  memcpy(&shmem_meta_obj_.shmem_->handle_, &ipc_handle, sizeof(T));
+  shmem_meta_obj_.shmem_->done_counter_ = 0;
+  return true;
+}
+
+template <typename T>
+bool MultiProcess<T>::WaitTillAllChildReads() {
+  size_t write_count = 0;
+  while (shmem_meta_obj_.shmem_->done_counter_ != num_proc_) {
+    ++write_count;
+  }
+  return true;
+}
+
+template <typename T>
+bool MultiProcess<T>::ReadHandleFromShmem(T& ipc_handle) {
+  size_t read_count = 0;
+  while (shmem_meta_obj_.shmem_->done_counter_ == -1) {
+    ++read_count;
+  }
+  memcpy(&ipc_handle, &shmem_meta_obj_.shmem_->handle_, sizeof(T));
+  return true;
+}
+
+template <typename T>
+bool MultiProcess<T>::NotifyParentDone() {
+  ++shmem_meta_obj_.shmem_->done_counter_;
+  return true;
+}
+
+#endif /* __unix__ */
diff --git a/projects/clr/hipamd/tests/src/ipc/hipMultiProcIpcEvent.cpp b/projects/clr/hipamd/tests/src/ipc/hipMultiProcIpcEvent.cpp
new file mode 100755
index 0000000000..ae1ce35a48
--- /dev/null
+++ b/projects/clr/hipamd/tests/src/ipc/hipMultiProcIpcEvent.cpp
@@ -0,0 +1,126 @@
+/*
+Copyright (c) 2015-2017 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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * TEST: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+#include "MultiProcess.h"
+
+void multi_process(int num_process, bool debug_process) {
+
+#ifdef __unix__
+
+  float *A_h, *B_h, *C_h;
+  float *A_d, *B_d, *C_d;
+  hipEvent_t start, stop;
+  size_t Nbytes = N * sizeof(float);
+
+  MultiProcess<hipIpcEventHandle_t>* mProcess = new MultiProcess<hipIpcEventHandle_t>(num_process);
+  mProcess->CreateShmem();
+  pid_t pid = mProcess->SpawnProcess(debug_process);
+
+  // Parent Process
+  if (pid != 0) {
+
+    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);
+
+    HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N);
+
+    // NULL stream check:
+    HIPCHECK(hipEventCreateWithFlags(&start, hipEventDisableTiming|hipEventInterprocess));
+    HIPCHECK(hipEventCreateWithFlags(&stop, hipEventDisableTiming|hipEventInterprocess));
+
+
+    HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+    HIPCHECK(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
+      HIPCHECK(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);
+
+
+      HIPCHECK(hipEventRecord(stop, NULL));
+      HIPCHECK(hipEventSynchronize(stop));
+      HIPCHECK(hipEventQuery(stop));
+      long long hostStop = HipTest::get_time();
+      //--- STOP TIMED REGION
+
+      float eventMs = 1.0f;
+      // should fail
+      HIPASSERT(hipSuccess != hipEventElapsedTime(&eventMs, start, stop));
+      float hostMs = HipTest::elapsed_time(hostStart, hostStop);
+
+      printf("host_time (gettimeofday)          =%6.3fms\n", hostMs);
+      printf("kernel_time (hipEventElapsedTime) =%6.3fms\n", eventMs);
+      printf("\n");
+
+    }
+
+    hipIpcEventHandle_t ipc_handle;
+    HIPCHECK(hipIpcGetEventHandle(&ipc_handle, start));
+
+    mProcess->WriteHandleToShmem(ipc_handle);
+    mProcess->WaitTillAllChildReads();
+
+  } else {
+    hipEvent_t ipc_event;
+    hipIpcEventHandle_t ipc_handle;
+    mProcess->ReadHandleFromShmem(ipc_handle);
+    HIPCHECK(hipIpcOpenEventHandle(&ipc_event, ipc_handle));
+
+    HIPCHECK(hipEventSynchronize(ipc_event));
+    HIPCHECK(hipEventDestroy(ipc_event));
+    mProcess->NotifyParentDone();
+  }
+
+  if (pid != 0) {
+    HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
+    printf("check:\n");
+    HipTest::checkVectorADD(A_h, B_h, C_h, N, true);
+
+    HIPCHECK(hipEventDestroy(start));
+    HIPCHECK(hipEventDestroy(stop));
+    delete mProcess;
+  }
+
+#endif /* __unix__ */
+
+}
+
+int main(int argc, char* argv[]) {
+  HipTest::parseStandardArguments(argc, argv, true);
+  multi_process((N < 64) ? N : 64, debug_test);
+  passed();
+}
diff --git a/projects/clr/hipamd/tests/src/ipc/hipMultiProcIpcMem.cpp b/projects/clr/hipamd/tests/src/ipc/hipMultiProcIpcMem.cpp
new file mode 100755
index 0000000000..15693bfbd6
--- /dev/null
+++ b/projects/clr/hipamd/tests/src/ipc/hipMultiProcIpcMem.cpp
@@ -0,0 +1,98 @@
+/*
+Copyright (c) 2015-2017 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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * TEST: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+#include "MultiProcess.h"
+
+#define NUM_ELEMS 1024
+#define OFFSET 128
+
+void multi_process(int num_process, bool debug_process) {
+
+#ifdef __unix__
+
+  int* ipc_dptr = nullptr;
+  int* ipc_hptr = nullptr;
+  int* ipc_out_dptr = nullptr;
+  int* ipc_out_hptr = nullptr;
+
+  MultiProcess<hipIpcMemHandle_t>* mProcess = new MultiProcess<hipIpcMemHandle_t>(num_process);
+  mProcess->CreateShmem();
+  pid_t pid = mProcess->SpawnProcess(debug_process);
+
+  // Parent Process
+  if (pid != 0) {
+    hipIpcMemHandle_t ipc_handle;
+    memset(&ipc_handle, 0x00, sizeof(hipIpcMemHandle_t));
+
+    HIPCHECK(hipMalloc((void**)&ipc_dptr, NUM_ELEMS * sizeof(int)));
+    HIPCHECK(hipIpcGetMemHandle(&ipc_handle, ipc_dptr));
+
+    ipc_hptr = new int[NUM_ELEMS];
+    for (size_t idx = 0; idx < NUM_ELEMS; ++idx) {
+      ipc_hptr[idx] = idx;
+    }
+
+    HIPCHECK(hipMemset(ipc_dptr, 0x00, (NUM_ELEMS * sizeof(int))));
+    HIPCHECK(hipMemcpy(ipc_dptr, ipc_hptr, (NUM_ELEMS * sizeof(int)), hipMemcpyHostToDevice));
+
+    mProcess->WriteHandleToShmem(ipc_handle);
+
+    mProcess->WaitTillAllChildReads();
+
+  } else {
+    ipc_out_hptr = new int[NUM_ELEMS];
+    memset(ipc_out_hptr, 0x00, (NUM_ELEMS * sizeof(int)));
+
+    hipIpcMemHandle_t ipc_handle;
+    mProcess->ReadHandleFromShmem(ipc_handle);
+    HIPCHECK(hipIpcOpenMemHandle((void**)&ipc_out_dptr, ipc_handle, 0));
+
+    HIPCHECK(hipMemcpy(ipc_out_hptr, ipc_out_dptr, (NUM_ELEMS * sizeof(int)),
+                       hipMemcpyDeviceToHost));
+    for (size_t idx = 0; idx < NUM_ELEMS; ++idx) {
+      if (ipc_out_hptr[idx] != idx) {
+        std::cout<<"Failing @ idx: "<< idx << std::endl;
+      }
+    }
+    mProcess->NotifyParentDone();
+    HIPCHECK(hipIpcCloseMemHandle(ipc_out_dptr));
+    delete[] ipc_out_hptr;
+  }
+
+  if (pid != 0) {
+    delete mProcess;
+  }
+
+#endif /* __unix__ */
+
+}
+
+
+int main(int argc, char* argv[]) {
+  HipTest::parseStandardArguments(argc, argv, true);
+  multi_process((N < 64) ? N : 64, debug_test);
+  passed();
+}
diff --git a/projects/clr/hipamd/tests/src/ipc/hipSimpleIpc.cpp b/projects/clr/hipamd/tests/src/ipc/hipSimpleIpcMem.cpp
similarity index 100%
rename from projects/clr/hipamd/tests/src/ipc/hipSimpleIpc.cpp
rename to projects/clr/hipamd/tests/src/ipc/hipSimpleIpcMem.cpp
diff --git a/projects/clr/hipamd/tests/src/test_common.cpp b/projects/clr/hipamd/tests/src/test_common.cpp
index f0a0883115..df4adc5e88 100644
--- a/projects/clr/hipamd/tests/src/test_common.cpp
+++ b/projects/clr/hipamd/tests/src/test_common.cpp
@@ -40,6 +40,7 @@ unsigned threadsPerBlock = 256;
 int p_gpuDevice = 0;
 unsigned p_verbose = 0;
 int p_tests = -1; /*which tests to run. Interpretation is left to each test.  default:all*/
+int debug_test = 0;
 #ifdef _WIN64
 const char* HIP_VISIBLE_DEVICES_STR = "HIP_VISIBLE_DEVICES=";
 const char* CUDA_VISIBLE_DEVICES_STR = "CUDA_VISIBLE_DEVICES=";
@@ -186,6 +187,10 @@ int parseStandardArguments(int argc, char* argv[], bool failOnUndefinedArg) {
                 failed("Bad tests argument");
             }
 
+        } else if (!strcmp(arg, "--debug") || (!strcmp(arg, "-d"))) {
+            if (++i >= argc || !HipTest::parseInt(argv[i], &debug_test)) {
+                failed("Bad tests argument");
+            }
         } else {
             if (failOnUndefinedArg) {
                 failed("Bad argument '%s'", arg);
diff --git a/projects/clr/hipamd/tests/src/test_common.h b/projects/clr/hipamd/tests/src/test_common.h
index e7612b079a..b6f03b1387 100755
--- a/projects/clr/hipamd/tests/src/test_common.h
+++ b/projects/clr/hipamd/tests/src/test_common.h
@@ -154,6 +154,7 @@ extern unsigned threadsPerBlock;
 extern int p_gpuDevice;
 extern unsigned p_verbose;
 extern int p_tests;
+extern int debug_test;
 extern const char* HIP_VISIBLE_DEVICES_STR;
 extern const char* CUDA_VISIBLE_DEVICES_STR;
 extern const char* PATH_SEPERATOR_STR;