SWDEV-276169 - Multiprocess IPC tests for Events and Memory.

Change-Id: I4a7af60e732de234a535574aa8597a7abd0b899b [ROCm/clr commit: 57615530a7]
2021-03-08 20:44:30 -05:00
@@ -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) {
@@ -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__ */
@@ -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();
+}
@@ -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();
+}
@@ -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);
@@ -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;