Added new ROCr Trap Handler Test

Intentionally trigger sw exceptions to verify GPU can handle abnormal scenarios

Change-Id: Ie80aa21883a912834ce6d917ecbb21ff6b3145f5
Signed-off-by: Chris Freehill <cfreehil@amd.com>


[ROCm/ROCR-Runtime commit: 6f6f02f352]

projects/rocr-runtime/rocrtst/suites/functional/trap_handler.cc

@@ -0,0 +1,381 @@
+/*
+ * =============================================================================
+ *   ROC Runtime Conformance Release License
+ * =============================================================================
+ * The University of Illinois/NCSA
+ * Open Source License (NCSA)
+ *
+ * Copyright (c) 2024, Advanced Micro Devices, Inc.
+ * All rights reserved.
+ *
+ * Developed by:
+ *
+ *                 AMD Research and AMD ROC Software Development
+ *
+ *                 Advanced Micro Devices, Inc.
+ *
+ *                 www.amd.com
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal with 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:
+ *
+ *  - Redistributions of source code must retain the above copyright notice,
+ *    this list of conditions and the following disclaimers.
+ *  - Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimers in
+ *    the documentation and/or other materials provided with the distribution.
+ *  - Neither the names of <Name of Development Group, Name of Institution>,
+ *    nor the names of its contributors may be used to endorse or promote
+ *    products derived from this Software without specific prior written
+ *    permission.
+ *
+ * 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
+ *
+ */
+
+#include "suites/functional/trap_handler.h"
+#include "common/base_rocr_utils.h"
+#include "common/common.h"
+#include "common/os.h"
+#include "common/helper_funcs.h"
+#include "gtest/gtest.h"
+#include "hsa/hsa.h"
+#include "hsa/hsa_ext_amd.h"
+
+#include <thread>
+
+#define RETRY_LIMIT 5
+#define DELAY_IN_MILLISECONDS 1
+#define TIMEOUT_LIMIT 5000  // milliseconds
+
+static const char kSubTestSeparator[] = "  **************************";
+static void PrintDebugSubtestHeader(const char* header) {
+  std::cout << "  *** TrapHandler Subtest: " << header << " ***" << std::endl;
+}
+
+static const uint32_t kNumBufferElements = 256;
+
+typedef struct q_callback_data_t {
+  hsa_queue_t** qptr;
+  bool trap;
+} q_callback_data;
+
+static void CallbackQueueErrorHandler(hsa_status_t status, hsa_queue_t* source, void* data);
+static hsa_status_t CallbackEventHandler(const hsa_amd_event_t* event, void* data);
+
+TrapHandler::TrapHandler(bool trigger_s_trap, bool trigger_memory_violation) : TestBase() {
+  std::string name;
+  std::string desc;
+
+  name = "ROCr Trap Handler Test";
+  desc =
+      "This set of tests intentionally trigger software exceptions and verify "
+      "that the GPU can handle abnormal situations.";
+
+  if (trigger_s_trap) {
+    name += ": Trigger a software trap";
+    desc +=
+        "\n\nCurrent sub-test intentionally triggers a software exception using"
+        " the 's_trap' instruction, to validate if the queue's error handling"
+        " callback is triggered.";
+  } else if (trigger_memory_violation) {
+    name += ": Trigger illegal memory access";
+    desc +=
+        "\n\nCurrent sub-test intentionally triggers a memory violation error"
+        " to attempt accessing an invalid memory address. It verifies if the "
+        " GPU Memory protection exception is triggered.";
+  }
+
+  set_title(name);
+  set_description(desc);
+  set_kernel_file_name("trap_handler_kernels.hsaco");
+  kernel_names_ = {"trigger_s_trap", "trigger_memory_violation"};
+}
+
+void TrapHandler::SetUp() {
+  hsa_status_t err;
+
+  TestBase::SetUp();
+
+  err = rocrtst::SetDefaultAgents(this);
+  CHECK(err);
+
+  err = rocrtst::SetPoolsTypical(this);
+  CHECK(err)
+
+  return;
+}
+
+static inline void AtomicSetPacketHeader(uint16_t header, uint16_t setup,
+                                         hsa_kernel_dispatch_packet_t* queue_packet) {
+  __atomic_store_n(reinterpret_cast<uint32_t*>(queue_packet), header | (setup << 16),
+                   __ATOMIC_RELEASE);
+}
+
+void TrapHandler::Run(void) {
+  if (!rocrtst::CheckProfile(this)) {
+    return;
+  }
+  TestBase::Run();
+}
+
+void TrapHandler::TriggerSoftwareTrap(void) {
+  std::vector<hsa_agent_t> cpus, gpus;
+  hsa_status_t err = hsa_iterate_agents(rocrtst::IterateCPUAgents, &cpus);
+  CHECK(err)
+
+  err = hsa_iterate_agents(rocrtst::IterateGPUAgents, &gpus);
+  CHECK(err)
+
+  // Check if there are any GPUs available
+  if (gpus.empty()) {
+    std::cerr << "No GPUs found." << std::endl;
+    return;
+  }
+
+  // Select the first GPU in the vector
+  hsa_agent_t selectedGpuAgent = gpus[0];
+  uint32_t node_id;
+  err = hsa_agent_get_info(selectedGpuAgent, HSA_AGENT_INFO_NODE, &node_id);
+  CHECK(err);
+
+  std::cout << "*** Running test on GPU node ID: 0x" << std::hex << node_id << "***\n" << std::endl;
+  execute_kernel("trigger_s_trap", cpus[0], selectedGpuAgent);
+  return;
+}
+
+void TrapHandler::TriggerMemoryViolation(void) {
+  std::vector<hsa_agent_t> cpus, gpus;
+  hsa_status_t err = hsa_iterate_agents(rocrtst::IterateCPUAgents, &cpus);
+  CHECK(err)
+
+  err = hsa_iterate_agents(rocrtst::IterateGPUAgents, &gpus);
+  CHECK(err)
+
+  // Check if there are any GPUs available
+  if (gpus.empty()) {
+    std::cerr << "No GPUs found." << std::endl;
+    return;
+  }
+
+  // Select the first GPU from the list
+  hsa_agent_t selectedGpuAgent = gpus[0];
+  uint32_t node_id;
+  err = hsa_agent_get_info(selectedGpuAgent, HSA_AGENT_INFO_NODE, &node_id);
+  CHECK(err);
+
+  std::cout << "*** Running test on GPU node ID: 0x" << std::hex << node_id << "***\n" << std::endl;
+  execute_kernel("trigger_memory_violation", cpus[0], selectedGpuAgent);
+  return;
+}
+
+void TrapHandler::execute_kernel(const char* kernel_name, hsa_agent_t cpuAgent,
+                                 hsa_agent_t gpuAgent) {
+  hsa_signal_t signal;
+  uint32_t queue_size = 0;
+
+  set_kernel_name(kernel_name);
+
+  hsa_status_t err = hsa_agent_get_info(gpuAgent, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_size);
+  CHECK(err)
+  q_callback_data data;
+  data.qptr = &queue;
+  data.trap = false;
+
+  // Create queue and register queue error handler callback function
+  err = hsa_queue_create(gpuAgent, queue_size, HSA_QUEUE_TYPE_MULTI, CallbackQueueErrorHandler,
+                         &data, 0, 0, &queue);
+  CHECK(err)
+  ASSERT_NE(queue, nullptr);
+  set_main_queue(queue);
+
+  // Register system event handler callback function
+  hsa_status_t status = hsa_amd_register_system_event_handler(&CallbackEventHandler, this);
+  ASSERT_EQ(status, HSA_STATUS_SUCCESS);
+
+  hsa_amd_memory_pool_t kernarg_pool;
+  err = hsa_amd_agent_iterate_memory_pools(cpuAgent, rocrtst::GetKernArgMemoryPool, &kernarg_pool);
+  CHECK(err)
+
+  hsa_amd_memory_pool_t global_pool;
+  err = hsa_amd_agent_iterate_memory_pools(cpuAgent, rocrtst::GetGlobalMemoryPool, &global_pool);
+  CHECK(err)
+
+  err = hsa_amd_memory_pool_allocate(global_pool, kNumBufferElements * sizeof(uint32_t), 0,
+                                     reinterpret_cast<void**>(&src_buffer_));
+
+  CHECK(err)
+
+  err = hsa_amd_agents_allow_access(1, &gpuAgent, NULL, src_buffer_);
+  CHECK(err)
+
+  for (uint32_t i = 0; i < kNumBufferElements; ++i) {
+    reinterpret_cast<uint32_t*>(src_buffer_)[i] = i;
+  }
+
+  err = hsa_amd_memory_pool_allocate(global_pool, kNumBufferElements * sizeof(uint32_t), 0,
+                                     reinterpret_cast<void**>(&dst_buffer_));
+  CHECK(err)
+  err = hsa_amd_agents_allow_access(1, &gpuAgent, NULL, dst_buffer_);
+  CHECK(err)
+
+  typedef struct __attribute__((aligned(16))) local_args_t {
+    uint32_t* dstArray;
+    uint32_t* srcArray;
+    uint32_t size;
+  } local_args;
+
+  local_args* KernArgs = NULL;
+  err = hsa_amd_memory_pool_allocate(kernarg_pool, sizeof(local_args), 0,
+                                     reinterpret_cast<void**>(&KernArgs));
+  CHECK(err)
+
+  err = hsa_amd_agents_allow_access(1, &gpuAgent, NULL, KernArgs);
+  CHECK(err)
+
+  KernArgs->dstArray = reinterpret_cast<uint32_t*>(dst_buffer_);
+  KernArgs->srcArray = reinterpret_cast<uint32_t*>(src_buffer_);
+  KernArgs->size = kNumBufferElements;
+
+  err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent);
+  CHECK(err)
+
+  err = hsa_signal_create(1, 0, NULL, &signal);
+  CHECK(err)
+
+  hsa_kernel_dispatch_packet_t aql;
+  memset(&aql, 0, sizeof(aql));
+
+  aql.header = 0;
+  aql.setup = 1;
+  aql.workgroup_size_x = kNumBufferElements;
+  aql.workgroup_size_y = 1;
+  aql.workgroup_size_z = 1;
+  aql.grid_size_x = kNumBufferElements;
+  aql.grid_size_y = 1;
+  aql.grid_size_z = 1;
+  aql.private_segment_size = 0;
+  aql.group_segment_size = 0;
+  aql.kernel_object = kernel_object();
+  aql.kernarg_address = KernArgs;
+  aql.completion_signal = signal;
+
+  const uint32_t queue_mask = queue->size - 1;
+
+  uint64_t index = hsa_queue_load_write_index_relaxed(queue);
+
+  hsa_queue_store_write_index_relaxed(queue, index + 1);
+
+  rocrtst::WriteAQLToQueueLoc(queue, index, &aql);
+
+  uint32_t aql_header = HSA_PACKET_TYPE_KERNEL_DISPATCH;
+
+  aql_header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE;
+  aql_header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;
+
+  void* q_base = queue->base_address;
+  rocrtst::AtomicSetPacketHeader(
+      aql_header, aql.setup,
+      &(reinterpret_cast<hsa_kernel_dispatch_packet_t*>(q_base))[index & queue_mask]);
+
+  hsa_signal_store_relaxed(queue->doorbell_signal, index);
+
+  hsa_signal_value_t completion = 0;
+  int retry_count = 0;
+
+  if (strcmp(kernel_name, "trigger_s_trap") == 0) {
+    completion = hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, 1, TIMEOUT_LIMIT,
+                                           HSA_WAIT_STATE_BLOCKED);
+    ASSERT_EQ(completion, 1);
+
+    while (data.trap != true && retry_count < RETRY_LIMIT) {
+      ++retry_count;
+      sleep(DELAY_IN_MILLISECONDS);
+    }
+    ASSERT_EQ(data.trap, true);
+  }
+
+  else if (strcmp(kernel_name, "trigger_memory_violation") == 0) {
+    while (event_occured != 1 && retry_count < RETRY_LIMIT) {
+      ++retry_count;
+      sleep(DELAY_IN_MILLISECONDS);
+    }
+    ASSERT_EQ(event_occured, 1);
+  }
+
+  // Cleanup
+  if (KernArgs) hsa_memory_free(KernArgs);
+  if (src_buffer_) hsa_memory_free(src_buffer_);
+  if (dst_buffer_) hsa_memory_free(dst_buffer_);
+  if (signal.handle) hsa_signal_destroy(signal);
+
+  if (!event_occured && queue) {
+    hsa_queue_destroy(queue);
+    queue = nullptr;
+  }
+
+  return;
+}
+
+void TrapHandler::DisplayTestInfo(void) { TestBase::DisplayTestInfo(); }
+
+void TrapHandler::DisplayResults(void) const {
+  if (!rocrtst::CheckProfile(this)) {
+    return;
+  }
+
+  TestBase::DisplayResults();
+  return;
+}
+
+void TrapHandler::Close() { TestBase::Close(); }
+
+hsa_status_t CallbackEventHandler(const hsa_amd_event_t* event, void* data) {
+  TrapHandler* ptr = reinterpret_cast<TrapHandler*>(data);
+
+  switch (event->event_type) {
+    case HSA_AMD_GPU_MEMORY_FAULT_EVENT:
+      ptr->event_occured = true;
+      printf("Subtest Passed: Runtime caught GPU Memory Fault Event successfully!\n");
+      break;
+
+    case HSA_AMD_GPU_HW_EXCEPTION_EVENT:
+      ptr->event_occured = true;
+      printf("Subtest Passed: Runtime caught GPU HW Exception Event successfully!\n");
+      break;
+    default:
+      // Unknown event type
+      printf("Subtest Failed: Unknown event type occurred\n");
+      break;
+  };
+
+  std::cout << kSubTestSeparator << std::endl;
+  // Returning success to indicate that event was handled
+  return HSA_STATUS_SUCCESS;
+}
+
+void CallbackQueueErrorHandler(hsa_status_t status, hsa_queue_t* source, void* data) {
+  std::cout << "Subtest Passed: Runtime caught trap instruction successfully!" << std::endl;
+
+  ASSERT_NE(source, nullptr);
+  ASSERT_NE(data, nullptr);
+
+  q_callback_data* debug_data = reinterpret_cast<q_callback_data*>(data);
+  hsa_queue_t* queue = *(debug_data->qptr);
+  //  check the queue id and user data
+  ASSERT_EQ(source->id, queue->id);
+  debug_data->trap = true;
+
+  std::cout << kSubTestSeparator << std::endl;
+  return;
+}

projects/rocr-runtime/rocrtst/suites/functional/trap_handler.h

@@ -0,0 +1,75 @@
+/*
+ * =============================================================================
+ *   ROC Runtime Conformance Release License
+ * =============================================================================
+ * The University of Illinois/NCSA
+ * Open Source License (NCSA)
+ *
+ * Copyright (c) 2024, Advanced Micro Devices, Inc.
+ * All rights reserved.
+ *
+ * Developed by:
+ *
+ *                 AMD Research and AMD ROC Software Development
+ *
+ *                 Advanced Micro Devices, Inc.
+ *
+ *                 www.amd.com
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal with 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:
+ *
+ *  - Redistributions of source code must retain the above copyright notice,
+ *    this list of conditions and the following disclaimers.
+ *  - Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimers in
+ *    the documentation and/or other materials provided with the distribution.
+ *  - Neither the names of <Name of Development Group, Name of Institution>,
+ *    nor the names of its contributors may be used to endorse or promote
+ *    products derived from this Software without specific prior written
+ *    permission.
+ *
+ * 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
+ *
+ */
+
+#ifndef ROCRTST_SUITES_FUNCTIONAL_TRAP_HANDLER_H_
+#define ROCRTST_SUITES_FUNCTIONAL_TRAP_HANDLER_H_
+
+#include "suites/test_common/test_base.h"
+#include "common/base_rocr.h"
+#include "common/common.h"
+
+class TrapHandler : public TestBase {
+ public:
+  bool event_occured = false;
+  hsa_queue_t* queue = nullptr;
+  std::vector<std::string> kernel_names_;
+
+  TrapHandler(bool trigger_s_trap, bool trigger_memory_violation);
+  virtual ~TrapHandler() {}
+  virtual void Run();
+  virtual void Close();
+  virtual void DisplayResults() const;
+  virtual void DisplayTestInfo(void);
+  void SetUp();
+  void TriggerSoftwareTrap(void);
+  void TriggerMemoryViolation(void);
+
+ private:
+  void* src_buffer_;
+  void* dst_buffer_;
+  void execute_kernel(const char* kernel_name, hsa_agent_t cpuAgent, hsa_agent_t gpuAgent);
+};
+
+#endif  // ROCRTST_SUITES_FUNCTIONAL_TRAP_HANDLER_H_

projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt

@@ -380,6 +380,11 @@ set(BITCODE_LIBS "${COMMON_BITCODE_LIBS}")
 set(CL_FILE_LIST "${KERNELS_DIR}/cu_mask_kernels.cl")
 build_sample_for_devices("cu_mask")
 
+#Trap Handler Test case
+set(BITCODE_LIBS "${COMMON_BITCODE_LIBS}")
+set(CL_FILE_LIST "${KERNELS_DIR}/trap_handler_kernels.cl")
+build_sample_for_devices("trap_handler")
+
 set(CMAKE_BUILD_WITH_INSTALL_RPATH ON)
 
 # Build rules

projects/rocr-runtime/rocrtst/suites/test_common/kernels/trap_handler_kernels.cl

@@ -0,0 +1,76 @@
+/*
+ * =============================================================================
+ *   ROC Runtime Conformance Release License
+ * =============================================================================
+ * The University of Illinois/NCSA
+ * Open Source License (NCSA)
+ *
+ * Copyright (c) 2024, Advanced Micro Devices, Inc.
+ * All rights reserved.
+ *
+ * Developed by:
+ *
+ *                 AMD Research and AMD ROC Software Development
+ *
+ *                 Advanced Micro Devices, Inc.
+ *
+ *                 www.amd.com
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal with 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:
+ *
+ *  - Redistributions of source code must retain the above copyright notice,
+ *    this list of conditions and the following disclaimers.
+ *  - Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimers in
+ *    the documentation and/or other materials provided with the distribution.
+ *  - Neither the names of <Name of Development Group, Name of Institution>,
+ *    nor the names of its contributors may be used to endorse or promote
+ *    products derived from this Software without specific prior written
+ *    permission.
+ *
+ * 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
+ *
+ */
+
+__kernel void trigger_s_trap(__global int* dstArray, __global const int* srcArray, const int sz) {
+  // Obtain the unique ID of the work-item
+  size_t id = get_global_id(0);
+
+  // Adding this line _just_ to avoid any compiler optimization of this kernel
+  dstArray[id] = srcArray[id] * srcArray[id];
+
+  // Explicitly trigger a software trap on the first thread using AMDGCN assembly
+  if (id == 0) {
+    __asm__ volatile("s_trap 1");
+  }
+}
+
+
+__kernel void trigger_memory_violation(__global int* dstArray, __global const int* srcArray,
+                                       const int sz) {
+  // Obtain the unique ID of the work-item
+  size_t id = get_global_id(0);
+
+  // Adding this line _just_ to avoid any compiler optimization of this kernel
+  dstArray[id] = srcArray[id] * srcArray[id];
+
+  if (id == 0) {
+    // Point to a garbage memory location
+    int* invalidPtr = (int*)0xdeadbeef;
+
+    // Try to read from and also write to the garbage memory location to trigger a memory violation
+    int invalidData = *invalidPtr;
+    *invalidPtr = 42;
+  }
+}

projects/rocr-runtime/rocrtst/suites/test_common/main.cc

@@ -77,6 +77,7 @@
 #include "suites/functional/aql_barrier_bit.h"
 #include "suites/functional/signal_kernel.h"
 #include "suites/functional/cu_masking.h"
+#include "suites/functional/trap_handler.h"
 #include "rocm_smi/rocm_smi.h"
 
 static RocrTstGlobals *sRocrtstGlvalues = nullptr;
@@ -220,6 +221,18 @@ TEST(rocrtstFunc, DISABLED_CU_Masking) {
   RunGenericTest(&sd);
 }
 
+TEST(rocrtstFunc, TriggerSoftwareTrap) {
+  TrapHandler trap(true, false);
+  RunCustomTestProlog(&trap);
+  trap.TriggerSoftwareTrap();
+}
+
+TEST(rocrtstFunc, TriggerMemoryViolation) {
+  TrapHandler trap(false, true);
+  RunCustomTestProlog(&trap);
+  trap.TriggerMemoryViolation();
+}
+
 #ifndef ROCRTST_EMULATOR_BUILD
 TEST(rocrtstFunc, IPC) {
   IPCTest ipc;