Added functional tests

[ROCm/rocshmem commit: f710a301fe]
2024-11-22 15:24:50 -06:00
commit fe7fb26f50
@@ -220,6 +220,33 @@ case $2 in
        echo "amoset_n2_w8_z1"
        ROC_SHMEM_MAX_NUM_CONTEXTS=8 mpirun -np 2 $1 -w 8 -z 1 -a 44 > $3/amoset_n2_w8_z1.log
        check amoset_n2_w8_z1
+        echo "putsignal_n2_w1_z1"
+        mpirun -np 2 $1 -w 1 -z 1 -s 1048576 -a 53 > $3/putsignal_n2_w1_z1.log
+        check putsignal_n2_w1_z1
+        echo "putsignalwg_n2_w2_z32"
+        ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 32 -s 1048576 -a 55 > $3/putsignalwg_n2_w2_z32.log
+        check putsignalwg_n2_w2_z32
+        echo "putsignalwave_n2_w1_z32"
+        mpirun -np 2 $1 -w 1 -z 32 -s 1048576 -a 55 > $3/putsignalwave_n2_w1_z32.log
+        check putsignalwave_n2_w1_z32
+        echo "putsignalnbi"
+        mpirun -np 2 $1 -w 1 -z 1 -s 1048576 -a 56 > $3/putsignalnbi_n2_w1_z1.log
+        check putsignalnbi_n2_w1_z1
+        echo "putsignalnbiwg"
+        ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 32 -s 1048576 -a 57 > $3/putsignalnbiwg_n2_w2_z32.log
+        check putsignalnbiwg_n2_w2_z32
+        echo "putsignalnbiwave"
+        mpirun -np 2 $1 -w 1 -z 32 -s 1048576 -a 58 > $3/putsignalnbiwave_n2_w1_z32.log
+        check putsignalnbiwave_n2_w1_z32
+        echo "signalfetch"
+        mpirun -np 2 $1 -w 1 -z 1 -a 59 > $3/signalfetch_n2_w1_z1.log
+        check signalfetch_n2_w1_z1
+        echo "signalfetchwg"
+        ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 32 -a 60 > $3/signalfetchwg_n2_w2_z32.log
+        check signalfetchwg_n2_w2_z32
+        echo "signalfetchwave"
+        mpirun -np 2 $1 -w 1 -z 32 -a 60 > $3/signalfetchwave_n2_w2_z32.log
+        check signalfetchwave_n2_w2_z32
        ;;

    ###########################################################################
@@ -576,6 +603,33 @@ case $2 in
    *"amoxor")
        mpirun -np 2 $1 -w 1 -z 1 -a 51
        ;;
+    *"putsignal")
+        mpirun -np 2 $1 -w 1 -z 1 -a 53
+        ;;
+    *"putsignalwg")
+        ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 32 -a 55
+        ;;
+    *"putsignalwave")
+        mpirun -np 2 $1 -w 1 -z 32 -a 55
+        ;;
+    *"putsignalnbi")
+        mpirun -np 2 $1 -w 1 -z 1 -a 56
+        ;;
+    *"putsignalnbiwg")
+        ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 32 -a 57
+        ;;
+    *"putsignalnbiwave")
+        mpirun -np 2 $1 -w 1 -z 32 -a 58
+        ;;
+    *"signalfetch")
+        mpirun -np 2 $1 -w 1 -z 1 -a 59
+        ;;
+    *"signalfetchwg")
+        ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 32 -a 60
+        ;;
+    *"signalfetchwave")
+        mpirun -np 2 $1 -w 1 -z 32 -a 60
+        ;;
    *)
        echo "UNKNOWN TEST TYPE: $2"
        exit -1
@@ -53,6 +53,8 @@ target_sources(
    swarm_tester.cpp
    random_access_tester.cpp
    shmem_ptr_tester.cpp
+    signaling_operations_tester.cpp
+    signaling_operations_tester.hpp
    extended_primitives.cpp
    empty_tester.cpp
    wave_level_primitives.cpp
@@ -0,0 +1,163 @@
+/******************************************************************************
+ * Copyright (c) 2024 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 "signaling_operations_tester.hpp"
+
+#include <roc_shmem/roc_shmem.hpp>
+
+using namespace rocshmem;
+
+/******************************************************************************
+ * DEVICE TEST KERNEL
+ *****************************************************************************/
+__global__ void SignalingOperationsTest(int loop, int skip, uint64_t *timer, char *s_buf,
+                              char *r_buf, int size, uint64_t *sig_addr,
+                              uint64_t *fetched_value,
+                              TestType type, ShmemContextType ctx_type) {
+  __shared__ roc_shmem_ctx_t ctx;
+  roc_shmem_wg_init();
+  roc_shmem_wg_ctx_create(ctx_type, &ctx);
+
+  uint64_t start;
+  uint64_t signal = 0;
+  int sig_op = ROC_SHMEM_SIGNAL_SET;
+
+  for (int i = 0; i < loop + skip; i++) {
+    if (i == skip) {
+        __syncthreads();
+        start = roc_shmem_timer();
+    }
+
+    switch (type) {
+      case PutSignalTestType:
+        roc_shmem_ctx_putmem_signal(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
+        break;
+      case WGPutSignalTestType:
+        roc_shmem_ctx_putmem_signal_wg(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
+        break;
+      case WAVEPutSignalTestType:
+        roc_shmem_ctx_putmem_signal_wave(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
+        break;
+      case PutSignalNBITestType:
+        roc_shmem_ctx_putmem_signal_nbi(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
+        break;
+      case WGPutSignalNBITestType:
+        roc_shmem_ctx_putmem_signal_nbi_wg(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
+        break;
+      case WAVEPutSignalNBITestType:
+        roc_shmem_ctx_putmem_signal_nbi_wave(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
+        break;
+      case SignalFetchTestType:
+        *fetched_value = roc_shmem_signal_fetch(sig_addr);
+        break;
+      case WGSignalFetchTestType:
+        *fetched_value = roc_shmem_signal_fetch_wg(sig_addr);
+        break;
+      case WAVESignalFetchTestType:
+        *fetched_value = roc_shmem_signal_fetch_wave(sig_addr);
+        break;
+      default:
+        break;
+    }
+  }
+
+  roc_shmem_ctx_quiet(ctx);
+
+  __syncthreads();
+
+  if (hipThreadIdx_x == 0) {
+    timer[hipBlockIdx_x] = roc_shmem_timer() - start;
+  }
+
+  roc_shmem_wg_ctx_destroy(&ctx);
+  roc_shmem_wg_finalize();
+}
+
+/******************************************************************************
+ * HOST TESTER CLASS METHODS
+ *****************************************************************************/
+SignalingOperationsTester::SignalingOperationsTester(TesterArguments args) : Tester(args) {
+  s_buf = (char *)roc_shmem_malloc(args.max_msg_size * args.wg_size);
+  r_buf = (char *)roc_shmem_malloc(args.max_msg_size * args.wg_size);
+  sig_addr = (uint64_t *)roc_shmem_malloc(sizeof(uint64_t));
+  CHECK_HIP(hipMallocManaged(&fetched_value, sizeof(uint64_t), hipMemAttachHost));
+}
+
+SignalingOperationsTester::~SignalingOperationsTester() {
+  roc_shmem_free(s_buf);
+  roc_shmem_free(r_buf);
+  roc_shmem_free(sig_addr);
+  CHECK_HIP(hipFree(fetched_value));
+}
+
+void SignalingOperationsTester::resetBuffers(uint64_t size) {
+  memset(s_buf, '0', args.max_msg_size * args.wg_size);
+  memset(r_buf, '1', args.max_msg_size * args.wg_size);
+  *fetched_value = -1;
+  *sig_addr = args.myid + 123;
+}
+
+void SignalingOperationsTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
+                                             uint64_t size) {
+  size_t shared_bytes = 0;
+
+  hipLaunchKernelGGL(SignalingOperationsTest, gridSize, blockSize, shared_bytes, stream,
+                     loop, args.skip, timer, s_buf, r_buf, size, sig_addr, fetched_value,
+                     _type, _shmem_context);
+
+  num_msgs = (loop + args.skip) * gridSize.x;
+  num_timed_msgs = loop;
+}
+
+void SignalingOperationsTester::verifyResults(uint64_t size) {
+  int check_data_id = (_type == PutSignalTestType ||
+                       _type == PutSignalNBITestType ||
+                       _type == WAVEPutSignalTestType ||
+                       _type == WAVEPutSignalNBITestType ||
+                       _type == WGPutSignalTestType ||
+                       _type == WGPutSignalNBITestType)
+                    ? 1 : -1; // do not check if it doesn't match a test
+
+  int check_fetched_value_id = (_type == SignalFetchTestType ||
+                                _type == WAVESignalFetchTestType ||
+                                _type == WGSignalFetchTestType)
+                             ? 0 : -1; // do not check if it doesn't match a test
+
+  if (args.myid == check_data_id) {
+    for (int i = 0; i < size; i++) {
+      if (r_buf[i] != '0') {
+        fprintf(stderr, "Data validation error at idx %d\n", i);
+        fprintf(stderr, "Got %c, Expected %c\n", r_buf[i], '0');
+        exit(-1);
+      }
+    }
+  }
+
+  if (args.myid == check_fetched_value_id) {
+    uint64_t value = *fetched_value;
+    uint64_t expected_value = (args.myid + 123);
+    if (value != expected_value) {
+      fprintf(stderr, "Fetched Value %lu, Expected %lu\n", value, expected_value);
+      exit(-1);
+    }
+  }
+}
@@ -0,0 +1,50 @@
+/******************************************************************************
+ * Copyright (c) 2024 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.
+ *****************************************************************************/
+
+#ifndef _SIGNALING_OPERATIONS_
+#define _SIGNALING_OPERATIONS_
+
+#include "tester.hpp"
+
+/******************************************************************************
+ * HOST TESTER CLASS
+ *****************************************************************************/
+class SignalingOperationsTester : public Tester {
+ public:
+  explicit SignalingOperationsTester(TesterArguments args);
+  virtual ~SignalingOperationsTester();
+
+ protected:
+  virtual void resetBuffers(uint64_t size) override;
+
+  virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
+                            uint64_t size) override;
+
+  virtual void verifyResults(uint64_t size) override;
+
+  char *s_buf = nullptr;
+  char *r_buf = nullptr;
+  uint64_t *sig_addr;
+  uint64_t *fetched_value;
+};
+
+#endif
@@ -44,6 +44,7 @@
 #include "primitive_tester.hpp"
 #include "random_access_tester.hpp"
 #include "shmem_ptr_tester.hpp"
+#include "signaling_operations_tester.hpp"
 #include "swarm_tester.hpp"
 #include "sync_tester.hpp"
 #include "team_broadcast_tester.hpp"
@@ -419,6 +420,42 @@ std::vector<Tester*> Tester::create(TesterArguments args) {
      }
      testers.push_back(new WaveLevelPrimitiveTester(args));
      return testers;
+    case PutSignalTestType:
+      if (rank == 0) std::cout << "Putmem Signal ###" << std::endl;
+      testers.push_back(new SignalingOperationsTester(args));
+      return testers;
+    case WGPutSignalTestType:
+      if (rank == 0) std::cout << "WG Putmem Signal ###" << std::endl;
+      testers.push_back(new SignalingOperationsTester(args));
+      return testers;
+    case WAVEPutSignalTestType:
+      if (rank == 0) std::cout << "Wave Putmem Signal ###" << std::endl;
+      testers.push_back(new SignalingOperationsTester(args));
+      return testers;
+    case PutSignalNBITestType:
+      if (rank == 0) std::cout << "Non-Blocking Putmem Signal ###" << std::endl;
+      testers.push_back(new SignalingOperationsTester(args));
+      return testers;
+    case WGPutSignalNBITestType:
+      if (rank == 0) std::cout << "Non-Blocking WG Putmem Signal ###" << std::endl;
+      testers.push_back(new SignalingOperationsTester(args));
+      return testers;
+    case WAVEPutSignalNBITestType:
+      if (rank == 0) std::cout << "Non-Blocking Wave Putmem Signal ###" << std::endl;
+      testers.push_back(new SignalingOperationsTester(args));
+      return testers;
+    case SignalFetchTestType:
+      if (rank == 0) std::cout << "Signal Fetch ###" << std::endl;
+      testers.push_back(new SignalingOperationsTester(args));
+      return testers;
+    case WGSignalFetchTestType:
+      if (rank == 0) std::cout << "WG Signal Fetch ###" << std::endl;
+      testers.push_back(new SignalingOperationsTester(args));
+      return testers;
+    case WAVESignalFetchTestType:
+      if (rank == 0) std::cout << "Wave Signal Fetch ###" << std::endl;
+      testers.push_back(new SignalingOperationsTester(args));
+      return testers;
    default:
      if (rank == 0) std::cout << "Unknown ###" << std::endl;
      testers.push_back(new PrimitiveTester(args));
@@ -84,7 +84,16 @@ enum TestType {
  AMO_AndTestType = 49,
  AMO_OrTestType = 50,
  AMO_XorTestType = 51,
-  PingAllTestType = 52
+  PingAllTestType = 52,
+  PutSignalTestType = 53,
+  WGPutSignalTestType = 54,
+  WAVEPutSignalTestType = 55,
+  PutSignalNBITestType = 56,
+  WGPutSignalNBITestType = 57,
+  WAVEPutSignalNBITestType = 58,
+  SignalFetchTestType = 59,
+  WGSignalFetchTestType = 60,
+  WAVESignalFetchTestType = 61,
 };

 enum OpType { PutType = 0, GetType = 1 };