From 2bbe21db560ae7d1e8688ce9a89175efe30a99f6 Mon Sep 17 00:00:00 2001
From: Aurelien Bouteiller <Aurelien.bouteiller@gmail.com>
Date: Fri, 2 May 2025 12:00:35 -0400
Subject: [PATCH] func-tests: Don't rely on asserts to catch invalid argv/env
 params (#96)

---
 .../functional_tests/random_access_tester.cpp |  9 ++++-
 .../team_ctx_infra_tester.cpp                 | 39 ++++++++++++++-----
 2 files changed, 37 insertions(+), 11 deletions(-)

diff --git a/tests/functional_tests/random_access_tester.cpp b/tests/functional_tests/random_access_tester.cpp
index 8c7d61b253..960046368d 100644
--- a/tests/functional_tests/random_access_tester.cpp
+++ b/tests/functional_tests/random_access_tester.cpp
@@ -136,7 +136,12 @@ RandomAccessTester::RandomAccessTester(TesterArguments args) : Tester(args) {
   int wg_size = args.wg_size;
   _num_waves = (args.wg_size / 64) * args.num_wgs;
   _num_bins = args.thread_access / args.coal_coef;
-  assert((args.wg_size / 64) <= 1);
+  if ((args.wg_size / 64) > 1 || (64 % _num_bins) != 0) {
+    printf("Argument are incorrect\n");
+    assert((args.wg_size / 64) <= 1);
+    assert((64 % _num_bins) == 0);
+    abort();
+  }
 
   s_buf = (int *)rocshmem_malloc(max_size * wg_size * space);
   r_buf = (int *)rocshmem_malloc(max_size * wg_size * space);
@@ -218,7 +223,7 @@ void RandomAccessTester::verifyResults(uint64_t size) {
   }
 
   CHECK_HIP(hipMemcpy(h_dev_buf, r_buf, space * args.wg_size * size,
-		      hipMemcpyDeviceToHost));
+                      hipMemcpyDeviceToHost));
   CHECK_HIP(hipDeviceSynchronize());
   for (uint64_t i = 0; i < (space * args.wg_size * size / sizeof(int)); i++) {
     if (h_dev_buf[i] != h_buf[i]) {
diff --git a/tests/functional_tests/team_ctx_infra_tester.cpp b/tests/functional_tests/team_ctx_infra_tester.cpp
index 1a973c55b8..73322eb109 100644
--- a/tests/functional_tests/team_ctx_infra_tester.cpp
+++ b/tests/functional_tests/team_ctx_infra_tester.cpp
@@ -51,17 +51,26 @@ __global__ void TeamCtxInfraTest(ShmemContextType ctx_type,
    */
 
   rocshmem_wg_team_create_ctx(team[0], ctx_type, &ctx1);
-  assert (nullptr != ctx1.ctx_opaque);
+  if (nullptr == ctx1.ctx_opaque) {
+    printf("Create ctx1 on team[0] returned an invalid context!\n");
+    abort();
+  }
   rocshmem_wg_team_create_ctx(team[0], ctx_type, &ctx2);
-  assert (nullptr != ctx2.ctx_opaque);
+  if (nullptr == ctx2.ctx_opaque) {
+    printf("Create ctx2 on team[0] returned an invalid context!\n");
+    abort();
+  }
   rocshmem_wg_ctx_destroy(&ctx1);
   rocshmem_wg_team_create_ctx(team[0], ctx_type, &ctx3);
-  assert (nullptr != ctx3.ctx_opaque);
+  if (nullptr == ctx3.ctx_opaque) {
+    printf("Create ctx3 on team[0] returned an invalid context!\n");
+    abort();
+  }
 
   __syncthreads();
 
   if (ctx3.team_opaque != ctx2.team_opaque) {
-    printf("Incorrect for teams of ctx2 and ctx3 to not equal each other\n");
+    printf("Incorrect for teams of ctx2 and ctx3 to be different!\n");
     abort();
   }
 
@@ -76,12 +85,14 @@ __global__ void TeamCtxInfraTest(ShmemContextType ctx_type,
    */
   for (int team_i = 0; team_i < NUM_TEAMS; team_i++) {
     rocshmem_wg_team_create_ctx(team[team_i], ctx_type, &ctx[team_i]);
-    assert (nullptr != ctx.ctx_opaque);
+    if (nullptr == ctx[team_i].ctx_opaque) {
+      printf("Create ctx on team[%d] returned an invalid context!\n", team_i);
+      abort();
+    }
   }
 
   if (ctx[0].team_opaque == ctx[NUM_TEAMS - 1].team_opaque) {
-    printf(
-        "Incorrect for ctx[0] team and ctx[NUM_TEAMS-1] to equal each other\n");
+    printf("Incorrect for teams of ctx[0] and ctx[NUM_TEAMS-1] to be equal to each other\n");
     abort();
   }
 
@@ -106,12 +117,22 @@ void TeamCtxInfraTester::resetBuffers(uint64_t size) {}
 void TeamCtxInfraTester::preLaunchKernel() {
   int n_pes = rocshmem_team_n_pes(ROCSHMEM_TEAM_WORLD);
 
+  // validate we can run the test
+  if (auto maximum_num_contexts_str = getenv("ROCSHMEM_MAX_NUM_CONTEXTS")) {
+    int max_ctx = atoi(maximum_num_contexts_str);
+    if (max_ctx <= NUM_TEAMS) {
+      printf("ROCSHMEM_MAX_NUM_CONTEXTS=%d is smaller than NUM_TEAMS %d, invalid test setup!\n", max_ctx, NUM_TEAMS);
+      assert(max_ctx > NUM_TEAMS);
+      abort();
+    }
+  }
+
   for (int team_i = 0; team_i < NUM_TEAMS; team_i++) {
     team_world_dup[team_i] = ROCSHMEM_TEAM_INVALID;
     rocshmem_team_split_strided(ROCSHMEM_TEAM_WORLD, 0, 1, n_pes, nullptr, 0,
                                  &team_world_dup[team_i]);
     if (team_world_dup[team_i] == ROCSHMEM_TEAM_INVALID) {
-      printf("Team %d is invalid!\n", team_i);
+      printf("Created team %d is invalid!\n", team_i);
       abort();
     }
   }
@@ -121,7 +142,7 @@ void TeamCtxInfraTester::preLaunchKernel() {
   rocshmem_team_split_strided(ROCSHMEM_TEAM_WORLD, 0, 1, n_pes, nullptr, 0,
                                &new_team);
   if (new_team != ROCSHMEM_TEAM_INVALID) {
-    printf("new team is not invalid\n");
+    printf("Created new team should have been invalid!\n");
     abort();
   }
 }