functional_tests: use size_t for size variable (#190)

Changed the data type of `size` to `size_t` in all functional tests to ensure
consistency with rocSHMEM APIs.

tests/functional_tests/amo_bitwise_tester.cpp

@@ -55,7 +55,7 @@ AMOBitwiseTester<T>::~AMOBitwiseTester() {
 }
 
 template <typename T>
-void AMOBitwiseTester<T>::resetBuffers(uint64_t size) {
+void AMOBitwiseTester<T>::resetBuffers(size_t size) {
   memset(_r_buf, 0, args.max_msg_size);
   memset(_ret_val, 0, args.max_msg_size * args.num_wgs);
   memset(_s_buf, 0, args.max_msg_size * args.num_wgs);
@@ -63,7 +63,7 @@ void AMOBitwiseTester<T>::resetBuffers(uint64_t size) {
 
 template <typename T>
 void AMOBitwiseTester<T>::launchKernel(dim3 gridsize, dim3 blocksize, int loop,
-                                       uint64_t size) {
+                                       size_t size) {
   size_t shared_bytes = 0;
 
   hipLaunchKernelGGL(AMOBitwiseTest, gridsize, blocksize, shared_bytes, stream,
@@ -82,7 +82,7 @@ void AMOBitwiseTester<T>::launchKernel(dim3 gridsize, dim3 blocksize, int loop,
 #endif
 
 template <typename T>
-void AMOBitwiseTester<T>::verifyResults(uint64_t size) {
+void AMOBitwiseTester<T>::verifyResults(size_t size) {
   T ret;
   if(DISABLE_IPC_TEST) {
     printf("AMO binary ops not implemented for IPC: values were not verified\n");

tests/functional_tests/amo_bitwise_tester.hpp

@@ -37,12 +37,12 @@ class AMOBitwiseTester : public Tester {
   virtual ~AMOBitwiseTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   dim3 _gridSize{};
   char *_r_buf;

tests/functional_tests/amo_extended_tester.cpp

@@ -55,7 +55,7 @@ AMOExtendedTester<T>::~AMOExtendedTester() {
 }
 
 template <typename T>
-void AMOExtendedTester<T>::resetBuffers(uint64_t size) {
+void AMOExtendedTester<T>::resetBuffers(size_t size) {
   memset(_r_buf, 0, args.max_msg_size);
   memset(_ret_val, 0, args.max_msg_size * args.num_wgs);
   memset(_s_buf, 0, args.max_msg_size * args.num_wgs);
@@ -63,7 +63,7 @@ void AMOExtendedTester<T>::resetBuffers(uint64_t size) {
 
 template <typename T>
 void AMOExtendedTester<T>::launchKernel(dim3 gridsize, dim3 blocksize, int loop,
-                                        uint64_t size) {
+                                        size_t size) {
   size_t shared_bytes = 0;
 
   hipLaunchKernelGGL(AMOExtendedTest, gridsize, blocksize, shared_bytes, stream,
@@ -82,7 +82,7 @@ void AMOExtendedTester<T>::launchKernel(dim3 gridsize, dim3 blocksize, int loop,
 #endif
 
 template <typename T>
-void AMOExtendedTester<T>::verifyResults(uint64_t size) {
+void AMOExtendedTester<T>::verifyResults(size_t size) {
   T ret;
   if (args.myid == 0) {
     T expected_val = 0;

tests/functional_tests/amo_extended_tester.hpp

@@ -37,12 +37,12 @@ class AMOExtendedTester : public Tester {
   virtual ~AMOExtendedTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   dim3 _gridSize{};
   char *_r_buf;

tests/functional_tests/amo_standard_tester.cpp

@@ -55,7 +55,7 @@ AMOStandardTester<T>::~AMOStandardTester() {
 }
 
 template <typename T>
-void AMOStandardTester<T>::resetBuffers(uint64_t size) {
+void AMOStandardTester<T>::resetBuffers(size_t size) {
   memset(_r_buf, 0, args.max_msg_size);
   memset(_ret_val, 0, args.max_msg_size * args.num_wgs);
   memset(_s_buf, 0, args.max_msg_size * args.num_wgs);
@@ -63,7 +63,7 @@ void AMOStandardTester<T>::resetBuffers(uint64_t size) {
 
 template <typename T>
 void AMOStandardTester<T>::launchKernel(dim3 gridsize, dim3 blocksize, int loop,
-                                        uint64_t size) {
+                                        size_t size) {
   size_t shared_bytes = 0;
 
   hipLaunchKernelGGL(AMOStandardTest, gridsize, blocksize, shared_bytes, stream,
@@ -76,7 +76,7 @@ void AMOStandardTester<T>::launchKernel(dim3 gridsize, dim3 blocksize, int loop,
 }
 
 template <typename T>
-void AMOStandardTester<T>::verifyResults(uint64_t size) {
+void AMOStandardTester<T>::verifyResults(size_t size) {
   T ret;
   if (args.myid == 0) {
     T expected_val = 0;

tests/functional_tests/amo_standard_tester.hpp

@@ -37,12 +37,12 @@ class AMOStandardTester : public Tester {
   virtual ~AMOStandardTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   dim3 _gridSize{};
   char *_r_buf;

tests/functional_tests/barrier_all_tester.cpp

@@ -95,7 +95,7 @@ BarrierAllTester::BarrierAllTester(TesterArguments args) : Tester(args) {}
 BarrierAllTester::~BarrierAllTester() {}
 
 void BarrierAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                                    uint64_t size) {
+                                    size_t size) {
   size_t shared_bytes = 0;
 
   hipLaunchKernelGGL(BarrierAllTest, gridSize, blockSize, shared_bytes, stream,
@@ -105,6 +105,6 @@ void BarrierAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
   num_timed_msgs = loop;
 }
 
-void BarrierAllTester::resetBuffers(uint64_t size) {}
+void BarrierAllTester::resetBuffers(size_t size) {}
 
-void BarrierAllTester::verifyResults(uint64_t size) {}
+void BarrierAllTester::verifyResults(size_t size) {}

tests/functional_tests/barrier_all_tester.hpp

@@ -41,12 +41,12 @@ class BarrierAllTester : public Tester {
   virtual ~BarrierAllTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 };
 
 #endif

tests/functional_tests/default_ctx_primitive_tester.cpp

@@ -34,7 +34,7 @@
  __global__ void DefaultCTXPrimitiveTest(int loop, int skip,
                                long long int *start_time,
                                long long int *end_time, char *source,
-                               char *dest, int size, TestType type,
+                               char *dest, size_t size, TestType type,
                                ShmemContextType ctx_type, int wf_size) {
    int wg_id = get_flat_grid_id();
    int t_id  = get_flat_block_id();
@@ -51,7 +51,7 @@
    /**
     * Calculate start index for each thread within the grid
     */
-   uint64_t offset = size * get_flat_id();
+   size_t offset = size * get_flat_id();
    source += offset;
    dest += offset;
  
@@ -155,13 +155,13 @@
    rocshmem_free(dest);
  }
  
- void DefaultCTXPrimitiveTester::resetBuffers(uint64_t size) {
+ void DefaultCTXPrimitiveTester::resetBuffers(size_t size) {
    size_t buff_size = size * args.wg_size * args.num_wgs;
    memset(dest, '1', buff_size);
  }
  
  void DefaultCTXPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize,
-                                              int loop, uint64_t size) {
+                                              int loop, size_t size) {
    size_t shared_bytes = 0;
  
    hipLaunchKernelGGL(DefaultCTXPrimitiveTest, gridSize, blockSize,
@@ -173,7 +173,7 @@
    num_timed_msgs = loop * gridSize.x * blockSize.x;
  }
  
- void DefaultCTXPrimitiveTester::verifyResults(uint64_t size) {
+ void DefaultCTXPrimitiveTester::verifyResults(size_t size) {
    int check_id =
        (_type == DefaultCTXGetTestType ||
         _type == DefaultCTXGetNBITestType || _type == DefaultCTXGTestType)
@@ -182,7 +182,7 @@
  
    if (args.myid == check_id) {
      size_t buff_size = size * args.wg_size * args.num_wgs;
-     for (uint64_t i = 0; i < buff_size; i++) {
+     for (size_t i = 0; i < buff_size; i++) {
        if (dest[i] != source[i]) {
          std::cerr << "Data validation error at idx " << i << std::endl;
          std::cerr << " Got " << dest[i] << ", Expected "

tests/functional_tests/default_ctx_primitive_tester.hpp

@@ -36,12 +36,12 @@ class DefaultCTXPrimitiveTester : public Tester {
   virtual ~DefaultCTXPrimitiveTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   char *source = nullptr;
   char *dest = nullptr;

tests/functional_tests/empty_tester.cpp

@@ -49,10 +49,10 @@ EmptyTester::EmptyTester(TesterArguments args) : Tester(args) {}
 
 EmptyTester::~EmptyTester() {}
 
-void EmptyTester::resetBuffers(uint64_t size) {}
+void EmptyTester::resetBuffers(size_t size) {}
 
 void EmptyTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                               uint64_t size) {
+                               size_t size) {
   size_t shared_bytes = 0;
 
   hipLaunchKernelGGL(EmptyTest, gridSize, blockSize, shared_bytes, stream,
@@ -60,4 +60,4 @@ void EmptyTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
                      _shmem_context);
 }
 
-void EmptyTester::verifyResults(uint64_t size) {}
+void EmptyTester::verifyResults(size_t size) {}

tests/functional_tests/empty_tester.hpp

@@ -36,12 +36,12 @@ class EmptyTester : public Tester {
   virtual ~EmptyTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 };
 
 #endif  // ROCSHMEM_CLIENTS_FUNCTIONAL_TESTS_EMPTY_TESTER_HPP

tests/functional_tests/ping_all_tester.cpp

@@ -77,13 +77,13 @@ PingAllTester::PingAllTester(TesterArguments args) : Tester(args) {
 
 PingAllTester::~PingAllTester() { rocshmem_free(r_buf); }
 
-void PingAllTester::resetBuffers(uint64_t size) {
+void PingAllTester::resetBuffers(size_t size) {
   int num_pes {rocshmem_n_pes()};
   memset(r_buf, 0, sizeof(int) * args.num_wgs * num_pes);
 }
 
 void PingAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                                  uint64_t size) {
+                                  size_t size) {
   size_t shared_bytes = 0;
 
   hipLaunchKernelGGL(PingAllTest, gridSize, blockSize, shared_bytes, stream,
@@ -94,4 +94,4 @@ void PingAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
   num_timed_msgs = loop;
 }
 
-void PingAllTester::verifyResults(uint64_t size) {}
+void PingAllTester::verifyResults(size_t size) {}

tests/functional_tests/ping_all_tester.hpp

@@ -42,12 +42,12 @@ class PingAllTester : public Tester {
   virtual ~PingAllTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   int *r_buf;
 };

tests/functional_tests/ping_pong_tester.cpp

@@ -77,12 +77,12 @@ PingPongTester::PingPongTester(TesterArguments args) : Tester(args) {
 
 PingPongTester::~PingPongTester() { rocshmem_free(r_buf); }
 
-void PingPongTester::resetBuffers(uint64_t size) {
+void PingPongTester::resetBuffers(size_t size) {
   memset(r_buf, 0, sizeof(int) * args.num_wgs);
 }
 
 void PingPongTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                                  uint64_t size) {
+                                  size_t size) {
   size_t shared_bytes = 0;
 
   hipLaunchKernelGGL(PingPongTest, gridSize, blockSize, shared_bytes, stream,
@@ -93,4 +93,4 @@ void PingPongTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
   num_timed_msgs = loop;
 }
 
-void PingPongTester::verifyResults(uint64_t size) {}
+void PingPongTester::verifyResults(size_t size) {}

tests/functional_tests/ping_pong_tester.hpp

@@ -42,12 +42,12 @@ class PingPongTester : public Tester {
   virtual ~PingPongTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   int *r_buf;
 };

tests/functional_tests/primitive_mr_tester.cpp

@@ -33,7 +33,7 @@ using namespace rocshmem;
  *****************************************************************************/
 __global__ void PrimitiveMRTest(int loop, long long int *start_time,
                                 long long int *end_time, char *s_buf,
-                                char *r_buf, int size,
+                                char *r_buf, size_t size,
                                 ShmemContextType ctx_type) {
   __shared__ rocshmem_ctx_t ctx;
   int wg_id = get_flat_grid_id();
@@ -79,7 +79,7 @@ void PrimitiveMRTester::resetBuffers(size_t size) {
 }
 
 void PrimitiveMRTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                                     uint64_t size) {
+                                     size_t size) {
   size_t shared_bytes = 0;
 
   /* Warmup */
@@ -98,7 +98,7 @@ void PrimitiveMRTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
   num_timed_msgs = loop * 64;
 }
 
-void PrimitiveMRTester::verifyResults(uint64_t size) {
+void PrimitiveMRTester::verifyResults(size_t size) {
   int check_id =
       (_type == GetTestType || _type == GetNBITestType || _type == GTestType)
           ? 0

tests/functional_tests/primitive_mr_tester.hpp

@@ -39,9 +39,9 @@ class PrimitiveMRTester : public Tester {
   virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   char *s_buf = nullptr;
   char *r_buf = nullptr;

tests/functional_tests/primitive_tester.cpp

@@ -33,7 +33,7 @@ using namespace rocshmem;
  *****************************************************************************/
 __global__ void PrimitiveTest(int loop, int skip, long long int *start_time,
                               long long int *end_time, char *source,
-                              char *dest, int size, TestType type,
+                              char *dest, size_t size, TestType type,
                               ShmemContextType ctx_type, int wf_size) {
   __shared__ rocshmem_ctx_t ctx;
   int wg_id = get_flat_grid_id();
@@ -52,7 +52,7 @@ __global__ void PrimitiveTest(int loop, int skip, long long int *start_time,
   /**
    * Calculate start index for each thread within the grid
    */
-  uint64_t offset = size * get_flat_id();
+  size_t offset = size * get_flat_id();
   source += offset;
   dest += offset;
 
@@ -156,13 +156,13 @@ PrimitiveTester::~PrimitiveTester() {
   rocshmem_free(dest);
 }
 
-void PrimitiveTester::resetBuffers(uint64_t size) {
+void PrimitiveTester::resetBuffers(size_t size) {
   size_t buff_size = size * args.wg_size * args.num_wgs;
   memset(dest, '1', buff_size);
 }
 
 void PrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                                   uint64_t size) {
+                                   size_t size) {
   size_t shared_bytes = 0;
 
   hipLaunchKernelGGL(PrimitiveTest, gridSize, blockSize, shared_bytes, stream,
@@ -173,7 +173,7 @@ void PrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
   num_timed_msgs = loop * gridSize.x * blockSize.x;
 }
 
-void PrimitiveTester::verifyResults(uint64_t size) {
+void PrimitiveTester::verifyResults(size_t size) {
   int check_id =
       (_type == GetTestType || _type == GetNBITestType || _type == GTestType)
           ? 0

tests/functional_tests/primitive_tester.hpp

@@ -36,12 +36,12 @@ class PrimitiveTester : public Tester {
   virtual ~PrimitiveTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   char *source = nullptr;
   char *dest = nullptr;

tests/functional_tests/random_access_tester.cpp

@@ -35,7 +35,8 @@ using namespace rocshmem;
 
 __device__ bool thread_passing(int num_bins, uint32_t *bin_threads,
                                uint32_t *off_bins, uint32_t *PE_bins,
-                               int *offset, int *PE, int coal_coef, int size) {
+                               size_t *offset, int *PE, int coal_coef,
+                               size_t size) {
   bool pass = false;
   int wave_id = ((hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x) /
                  64);  // get_global_wave_id();
@@ -55,7 +56,7 @@ __device__ bool thread_passing(int num_bins, uint32_t *bin_threads,
 
 __global__ void RandomAccessTest(int loop, int skip, long long int *start_time,
                                  long long int *end_time, int *s_buf,
-                                 int *r_buf, int size, OpType type,
+                                 int *r_buf, size_t size, OpType type,
                                  int coal_coef, int num_bins, int num_waves,
                                  uint32_t *threads_bins, uint32_t *off_bins,
                                  uint32_t *PE_bins, ShmemContextType ctx_type) {
@@ -65,7 +66,7 @@ __global__ void RandomAccessTest(int loop, int skip, long long int *start_time,
   rocshmem_wg_ctx_create(ctx_type, &ctx);
 
   int pe = rocshmem_ctx_my_pe(ctx);
-  int offset;
+  size_t offset;
   int PE;
 
   if (thread_passing(num_bins, threads_bins, off_bins, PE_bins, &offset, &PE,
@@ -166,7 +167,7 @@ RandomAccessTester::~RandomAccessTester() {
   CHECK_HIP(hipFree(_PE_bins));
 }
 
-void RandomAccessTester::resetBuffers(uint64_t size) {
+void RandomAccessTester::resetBuffers(size_t size) {
   for (size_t i = 0; i < args.max_msg_size / sizeof(int) * args.wg_size * space;
        i++) {
     s_buf[i] = 1;
@@ -176,7 +177,7 @@ void RandomAccessTester::resetBuffers(uint64_t size) {
 }
 
 void RandomAccessTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                                      uint64_t size) {
+                                      size_t size) {
   size_t shared_bytes = 0;
 
   int _thread_access = args.thread_access;
@@ -200,7 +201,7 @@ void RandomAccessTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
   num_timed_msgs = loop * _num_waves * _thread_access;
 }
 
-void RandomAccessTester::verifyResults(uint64_t size) {
+void RandomAccessTester::verifyResults(size_t size) {
   uint64_t offset;
   for (int k = 0; k < _num_waves; k++) {
     for (int i = 0; i < _num_bins; i++) {

tests/functional_tests/random_access_tester.hpp

@@ -46,12 +46,12 @@ class RandomAccessTester : public Tester {
   virtual ~RandomAccessTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   int *r_buf;
   int *s_buf;

tests/functional_tests/shmem_ptr_tester.cpp

@@ -59,13 +59,13 @@ ShmemPtrTester::~ShmemPtrTester() {
   rocshmem_free(r_buf);
 }
 
-void ShmemPtrTester::resetBuffers(uint64_t size) {
+void ShmemPtrTester::resetBuffers(size_t size) {
   memset(r_buf, '0', args.max_msg_size);
   memset(_available, 0, sizeof(int));
 }
 
 void ShmemPtrTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                                  uint64_t size) {
+                                  size_t size) {
   size_t shared_bytes = 0;
 
   hipLaunchKernelGGL(ShmemPtrTest, gridSize, blockSize, shared_bytes, stream,
@@ -75,7 +75,7 @@ void ShmemPtrTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
   num_timed_msgs = 0;
 }
 
-void ShmemPtrTester::verifyResults(uint64_t size) {
+void ShmemPtrTester::verifyResults(size_t size) {
   if (args.myid == 0) {
     if (*_available == 0) {
       fprintf(stderr, "SHMEM_PTR NOT AVAILBLE \n");

tests/functional_tests/shmem_ptr_tester.hpp

@@ -36,12 +36,12 @@ class ShmemPtrTester : public Tester {
   virtual ~ShmemPtrTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   char *r_buf = nullptr;
   int *_available = nullptr;

tests/functional_tests/signaling_operations_tester.cpp

@@ -33,8 +33,9 @@ using namespace rocshmem;
  *****************************************************************************/
 __global__ void PutmemSignalTest(int loop, int skip, long long int *start_time,
                                  long long int *end_time, char *s_buf,
-                                 char *r_buf, int size, uint64_t *sig_addr,
-                                 TestType type, ShmemContextType ctx_type, int sig_op) {
+                                 char *r_buf, size_t size, uint64_t *sig_addr,
+                                 TestType type, ShmemContextType ctx_type,
+                                 int sig_op) {
   __shared__ rocshmem_ctx_t ctx;
   int wg_id = get_flat_grid_id();
   rocshmem_wg_init();
@@ -50,22 +51,28 @@ __global__ void PutmemSignalTest(int loop, int skip, long long int *start_time,
 
     switch (type) {
       case PutSignalTestType:
-        rocshmem_ctx_putmem_signal(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
+        rocshmem_ctx_putmem_signal(ctx, r_buf, s_buf, size, sig_addr,
+                                   signal, sig_op, 1);
         break;
       case WGPutSignalTestType:
-        rocshmem_ctx_putmem_signal_wg(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
+        rocshmem_ctx_putmem_signal_wg(ctx, r_buf, s_buf, size, sig_addr,
+                                      signal, sig_op, 1);
         break;
       case WAVEPutSignalTestType:
-        rocshmem_ctx_putmem_signal_wave(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
+        rocshmem_ctx_putmem_signal_wave(ctx, r_buf, s_buf, size, sig_addr,
+                                        signal, sig_op, 1);
         break;
       case PutSignalNBITestType:
-        rocshmem_ctx_putmem_signal_nbi(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
+        rocshmem_ctx_putmem_signal_nbi(ctx, r_buf, s_buf, size, sig_addr,
+                                       signal, sig_op, 1);
         break;
       case WGPutSignalNBITestType:
-        rocshmem_ctx_putmem_signal_nbi_wg(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
+        rocshmem_ctx_putmem_signal_nbi_wg(ctx, r_buf, s_buf, size, sig_addr,
+                                          signal, sig_op, 1);
         break;
       case WAVEPutSignalNBITestType:
-        rocshmem_ctx_putmem_signal_nbi_wave(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
+        rocshmem_ctx_putmem_signal_nbi_wave(ctx, r_buf, s_buf, size, sig_addr,
+                                            signal, sig_op, 1);
         break;
       default:
         break;
@@ -124,15 +131,17 @@ __global__ void SignalFetchTest(int loop, int skip, long long int *start_time,
 /******************************************************************************
  * HOST TESTER CLASS METHODS
  *****************************************************************************/
-SignalingOperationsTester::SignalingOperationsTester(TesterArguments args) : Tester(args) {
+SignalingOperationsTester::SignalingOperationsTester(TesterArguments args)
+  : Tester(args) {
   s_buf = (char *)rocshmem_malloc(args.max_msg_size * args.wg_size);
   r_buf = (char *)rocshmem_malloc(args.max_msg_size * args.wg_size);
   sig_addr = (uint64_t *)rocshmem_malloc(sizeof(uint64_t));
   CHECK_HIP(hipMallocManaged(&fetched_value, sizeof(uint64_t), hipMemAttachHost));
 }
 
-SignalingOperationsTester::SignalingOperationsTester(TesterArguments args, int signal_op)
-                          : SignalingOperationsTester(args) {
+SignalingOperationsTester::SignalingOperationsTester(TesterArguments args,
+                                                     int signal_op)
+  : SignalingOperationsTester(args) {
   sig_op = signal_op;
 }
 
@@ -143,7 +152,7 @@ SignalingOperationsTester::~SignalingOperationsTester() {
   CHECK_HIP(hipFree(fetched_value));
 }
 
-void SignalingOperationsTester::resetBuffers(uint64_t size) {
+void SignalingOperationsTester::resetBuffers(size_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;
@@ -151,7 +160,7 @@ void SignalingOperationsTester::resetBuffers(uint64_t size) {
 }
 
 void SignalingOperationsTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                                             uint64_t size) {
+                                             size_t size) {
   size_t shared_bytes = 0;
 
 
@@ -170,7 +179,7 @@ void SignalingOperationsTester::launchKernel(dim3 gridSize, dim3 blockSize, int
   num_timed_msgs = loop;
 }
 
-void SignalingOperationsTester::verifyResults(uint64_t size) {
+void SignalingOperationsTester::verifyResults(size_t size) {
   if (_type == SignalFetchTestType     ||
       _type == WAVESignalFetchTestType ||
       _type == WGSignalFetchTestType) {

tests/functional_tests/signaling_operations_tester.hpp

@@ -37,12 +37,12 @@ class SignalingOperationsTester : public Tester {
   virtual ~SignalingOperationsTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   int sig_op;
   char *s_buf = nullptr;

tests/functional_tests/sync_all_tester.cpp

@@ -95,7 +95,7 @@ SyncAllTester::SyncAllTester(TesterArguments args) : Tester(args) {}
 SyncAllTester::~SyncAllTester() {}
 
 void SyncAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                                    uint64_t size) {
+                                    size_t size) {
   size_t shared_bytes = 0;
 
   hipLaunchKernelGGL(SyncAllTest, gridSize, blockSize, shared_bytes, stream,
@@ -105,6 +105,6 @@ void SyncAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
   num_timed_msgs = loop;
 }
 
-void SyncAllTester::resetBuffers(uint64_t size) {}
+void SyncAllTester::resetBuffers(size_t size) {}
 
-void SyncAllTester::verifyResults(uint64_t size) {}
+void SyncAllTester::verifyResults(size_t size) {}

tests/functional_tests/sync_all_tester.hpp

@@ -41,12 +41,12 @@ class SyncAllTester : public Tester {
   virtual ~SyncAllTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 };
 
 #endif

tests/functional_tests/sync_tester.cpp

@@ -90,7 +90,7 @@ SyncTester::~SyncTester() {
   CHECK_HIP(hipFree(team_sync_world_dup));
 }
 
-void SyncTester::resetBuffers(uint64_t size) {}
+void SyncTester::resetBuffers(size_t size) {}
 
 void SyncTester::preLaunchKernel() {
   int n_pes = rocshmem_team_n_pes(ROCSHMEM_TEAM_WORLD);
@@ -107,7 +107,7 @@ void SyncTester::preLaunchKernel() {
 }
 
 void SyncTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                              uint64_t size) {
+                              size_t size) {
   size_t shared_bytes = 0;
 
   int n_pes = rocshmem_team_n_pes(ROCSHMEM_TEAM_WORLD);
@@ -126,4 +126,4 @@ void SyncTester::postLaunchKernel() {
   }
 }
 
-void SyncTester::verifyResults(uint64_t size) {}
+void SyncTester::verifyResults(size_t size) {}

tests/functional_tests/sync_tester.hpp

@@ -40,16 +40,16 @@ class SyncTester : public Tester {
   virtual ~SyncTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void preLaunchKernel() override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
   virtual void postLaunchKernel() override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
  private:
   /**

tests/functional_tests/team_alltoall_tester.cpp

@@ -153,7 +153,7 @@ void TeamAlltoallTester<T1>::preLaunchKernel() {
 
 template <typename T1>
 void TeamAlltoallTester<T1>::launchKernel(dim3 gridSize, dim3 blockSize,
-                                          int loop, uint64_t size) {
+                                          int loop, size_t size) {
   size_t shared_bytes = 0;
 
   int num_elems = size / sizeof(T1);
@@ -175,7 +175,7 @@ void TeamAlltoallTester<T1>::postLaunchKernel() {
 }
 
 template <typename T1>
-void TeamAlltoallTester<T1>::resetBuffers(uint64_t size) {
+void TeamAlltoallTester<T1>::resetBuffers(size_t size) {
 
   int num_elems = size / sizeof(T1);
   int buff_size = num_elems * sizeof(T1) * args.num_wgs * n_pes;
@@ -204,7 +204,7 @@ void TeamAlltoallTester<T1>::resetBuffers(uint64_t size) {
 }
 
 template <typename T1>
-void TeamAlltoallTester<T1>::verifyResults(uint64_t size) {
+void TeamAlltoallTester<T1>::verifyResults(size_t size) {
   int num_elems = size / sizeof(T1);
   int idx = 0;
 

tests/functional_tests/team_alltoall_tester.hpp

@@ -42,16 +42,16 @@ class TeamAlltoallTester : public Tester {
   virtual ~TeamAlltoallTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void preLaunchKernel() override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
   virtual void postLaunchKernel() override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   T1 *source_buf = nullptr;
   T1 *dest_buf = nullptr;

tests/functional_tests/team_barrier_tester.cpp

@@ -106,7 +106,7 @@ void TeamBarrierTester::preLaunchKernel() {
 }
 
 void TeamBarrierTester::launchKernel(dim3 gridSize, dim3 blockSize,
-                                           int loop, uint64_t size) {
+                                           int loop, size_t size) {
   size_t shared_bytes = 0;
 
   hipLaunchKernelGGL(TeamBarrierTest, gridSize, blockSize, shared_bytes,
@@ -124,6 +124,6 @@ void TeamBarrierTester::postLaunchKernel() {
   }
 }
 
-void TeamBarrierTester::resetBuffers(uint64_t size) {}
+void TeamBarrierTester::resetBuffers(size_t size) {}
 
-void TeamBarrierTester::verifyResults(uint64_t size) {}
+void TeamBarrierTester::verifyResults(size_t size) {}

tests/functional_tests/team_barrier_tester.hpp

@@ -41,16 +41,16 @@ class TeamBarrierTester : public Tester {
   virtual ~TeamBarrierTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void preLaunchKernel() override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
   virtual void postLaunchKernel() override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
  private:
   int my_pe = 0;

tests/functional_tests/team_broadcast_tester.cpp

@@ -153,7 +153,7 @@ void TeamBroadcastTester<T1>::preLaunchKernel() {
 
 template <typename T1>
 void TeamBroadcastTester<T1>::launchKernel(dim3 gridSize, dim3 blockSize,
-                                           int loop, uint64_t size) {
+                                           int loop, size_t size) {
   size_t shared_bytes = 0;
 
   int num_elems = size / sizeof(T1);
@@ -175,7 +175,7 @@ void TeamBroadcastTester<T1>::postLaunchKernel() {
 }
 
 template <typename T1>
-void TeamBroadcastTester<T1>::resetBuffers(uint64_t size) {
+void TeamBroadcastTester<T1>::resetBuffers(size_t size) {
 
   int num_elems = size / sizeof(T1);
   int buff_size = num_elems * sizeof(T1) * args.num_wgs;
@@ -203,7 +203,7 @@ void TeamBroadcastTester<T1>::resetBuffers(uint64_t size) {
 }
 
 template <typename T1>
-void TeamBroadcastTester<T1>::verifyResults(uint64_t size) {
+void TeamBroadcastTester<T1>::verifyResults(size_t size) {
 
   int num_elems = size / sizeof(T1);
   int idx = 0;

tests/functional_tests/team_broadcast_tester.hpp

@@ -42,16 +42,16 @@ class TeamBroadcastTester : public Tester {
   virtual ~TeamBroadcastTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void preLaunchKernel() override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
   virtual void postLaunchKernel() override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   T1 *source_buf;
   T1 *dest_buf;

tests/functional_tests/team_ctx_infra_tester.cpp

@@ -113,7 +113,7 @@ TeamCtxInfraTester::TeamCtxInfraTester(TesterArguments args) : Tester(args) {}
 
 TeamCtxInfraTester::~TeamCtxInfraTester() {}
 
-void TeamCtxInfraTester::resetBuffers(uint64_t size) {}
+void TeamCtxInfraTester::resetBuffers(size_t size) {}
 
 void TeamCtxInfraTester::preLaunchKernel() {
   int n_pes = rocshmem_team_n_pes(ROCSHMEM_TEAM_WORLD);
@@ -149,7 +149,7 @@ void TeamCtxInfraTester::preLaunchKernel() {
 }
 
 void TeamCtxInfraTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                                      uint64_t size) {
+                                      size_t size) {
   size_t shared_bytes = 0;
 
   /* Copy array of teams to device */
@@ -170,4 +170,4 @@ void TeamCtxInfraTester::postLaunchKernel() {
   }
 }
 
-void TeamCtxInfraTester::verifyResults(uint64_t size) {}
+void TeamCtxInfraTester::verifyResults(size_t size) {}

tests/functional_tests/team_ctx_infra_tester.hpp

@@ -36,16 +36,16 @@ class TeamCtxInfraTester : public Tester {
   virtual ~TeamCtxInfraTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void preLaunchKernel() override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
   virtual void postLaunchKernel() override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   char *s_buf = nullptr;
   char *r_buf = nullptr;

tests/functional_tests/team_ctx_primitive_tester.cpp

@@ -35,7 +35,7 @@ rocshmem_team_t team_primitive_world_dup;
  *****************************************************************************/
 __global__ void TeamCtxPrimitiveTest(int loop, int skip, long long int *start_time,
                                      long long int *end_time, char *source,
-                                     char *dest, int size, TestType type,
+                                     char *dest, size_t size, TestType type,
                                      ShmemContextType ctx_type, int wf_size,
                                      rocshmem_team_t team) {
   __shared__ rocshmem_ctx_t ctx;
@@ -56,7 +56,7 @@ __global__ void TeamCtxPrimitiveTest(int loop, int skip, long long int *start_ti
   /**
    * Calculate start index for each thread within the grid
    */
-  uint64_t offset = size * get_flat_id();
+  size_t offset = size * get_flat_id();
   source += offset;
   dest += offset;
 
@@ -148,7 +148,7 @@ TeamCtxPrimitiveTester::~TeamCtxPrimitiveTester() {
   rocshmem_free(dest);
 }
 
-void TeamCtxPrimitiveTester::resetBuffers(uint64_t size) {
+void TeamCtxPrimitiveTester::resetBuffers(size_t size) {
   size_t buff_size = size * args.wg_size * args.num_wgs;
   memset(dest, '1', buff_size);
 }
@@ -162,7 +162,7 @@ void TeamCtxPrimitiveTester::preLaunchKernel() {
 }
 
 void TeamCtxPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize,
-                                          int loop, uint64_t size) {
+                                          int loop, size_t size) {
   size_t shared_bytes = 0;
 
   hipLaunchKernelGGL(TeamCtxPrimitiveTest, gridSize, blockSize, shared_bytes,
@@ -178,7 +178,7 @@ void TeamCtxPrimitiveTester::postLaunchKernel() {
   rocshmem_team_destroy(team_primitive_world_dup);
 }
 
-void TeamCtxPrimitiveTester::verifyResults(uint64_t size) {
+void TeamCtxPrimitiveTester::verifyResults(size_t size) {
   int check_id =
       (_type == TeamCtxGetTestType || _type == TeamCtxGetNBITestType) ? 0 : 1;
 

tests/functional_tests/team_ctx_primitive_tester.hpp

@@ -36,16 +36,16 @@ class TeamCtxPrimitiveTester : public Tester {
   virtual ~TeamCtxPrimitiveTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void preLaunchKernel() override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
   virtual void postLaunchKernel() override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   char *source = nullptr;
   char *dest = nullptr;

tests/functional_tests/team_fcollect_tester.cpp

@@ -168,7 +168,7 @@ void TeamFcollectTester<T1>::preLaunchKernel() {
 
 template <typename T1>
 void TeamFcollectTester<T1>::launchKernel(dim3 gridSize, dim3 blockSize,
-                                          int loop, uint64_t size) {
+                                          int loop, size_t size) {
   size_t shared_bytes = 0;
 
   int num_elems = size / sizeof(T1);
@@ -193,7 +193,7 @@ void TeamFcollectTester<T1>::postLaunchKernel() {
 }
 
 template <typename T1>
-void TeamFcollectTester<T1>::resetBuffers(uint64_t size) {
+void TeamFcollectTester<T1>::resetBuffers(size_t size) {
   int num_elems = (size / sizeof(T1));
   int buff_size = num_elems * sizeof(T1) * args.num_wgs * n_pes;
 
@@ -201,7 +201,7 @@ void TeamFcollectTester<T1>::resetBuffers(uint64_t size) {
 }
 
 template <typename T1>
-void TeamFcollectTester<T1>::verifyResults(uint64_t size) {
+void TeamFcollectTester<T1>::verifyResults(size_t size) {
 
   int num_elems = size / sizeof(T1);
   int idx = 0;

tests/functional_tests/team_fcollect_tester.hpp

@@ -42,16 +42,16 @@ class TeamFcollectTester : public Tester {
   virtual ~TeamFcollectTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void preLaunchKernel() override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
   virtual void postLaunchKernel() override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   T1 *source_buf;
   T1 *dest_buf;

tests/functional_tests/team_reduction_tester.cpp

@@ -77,7 +77,7 @@ rocshmem_team_t team_reduce_world_dup;
 template <typename T1, ROCSHMEM_OP T2>
 __global__ void TeamReductionTest(int loop, int skip, long long int *start_time,
                                   long long int *end_time, T1 *s_buf, T1 *r_buf,
-                                  int size, TestType type,
+                                  size_t size, TestType type,
                                   ShmemContextType ctx_type,
                                   rocshmem_team_t team) {
   __shared__ rocshmem_ctx_t ctx;
@@ -136,7 +136,7 @@ void TeamReductionTester<T1, T2>::preLaunchKernel() {
 
 template <typename T1, ROCSHMEM_OP T2>
 void TeamReductionTester<T1, T2>::launchKernel(dim3 gridSize, dim3 blockSize,
-                                               int loop, uint64_t size) {
+                                               int loop, size_t size) {
   size_t shared_bytes = 0;
 
   hipLaunchKernelGGL(HIP_KERNEL_NAME(TeamReductionTest<T1, T2>), gridSize,
@@ -154,14 +154,14 @@ void TeamReductionTester<T1, T2>::postLaunchKernel() {
 }
 
 template <typename T1, ROCSHMEM_OP T2>
-void TeamReductionTester<T1, T2>::resetBuffers(uint64_t size) {
+void TeamReductionTester<T1, T2>::resetBuffers(size_t size) {
   for (uint64_t i = 0; i < args.max_msg_size; i++) {
     init_buf(s_buf[i], r_buf[i]);
   }
 }
 
 template <typename T1, ROCSHMEM_OP T2>
-void TeamReductionTester<T1, T2>::verifyResults(uint64_t size) {
+void TeamReductionTester<T1, T2>::verifyResults(size_t size) {
   int n_pes = rocshmem_n_pes();
   for (uint64_t i = 0; i < size; i++) {
     auto r = verify_buf(r_buf[i], (T1)n_pes);

tests/functional_tests/team_reduction_tester.hpp

@@ -42,16 +42,16 @@ class TeamReductionTester : public Tester {
   virtual ~TeamReductionTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void preLaunchKernel() override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
   virtual void postLaunchKernel() override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   T1 *s_buf;
   T1 *r_buf;

tests/functional_tests/tester.cpp

@@ -478,7 +478,7 @@ void Tester::execute() {
    * Some tests loop through data sizes in powers of 2 and report the
    * results for those ranges.
    */
-  for (uint64_t size = args.min_msg_size; size <= args.max_msg_size;
+  for (size_t size = args.min_msg_size; size <= args.max_msg_size;
        size <<= 1) {
     resetBuffers(size);
 

tests/functional_tests/tester_arguments.hpp

@@ -54,8 +54,8 @@ class TesterArguments {
   unsigned num_wgs = 1;
   unsigned num_threads = 1;
   unsigned algorithm = 0;
-  uint64_t min_msg_size = 1;
-  uint64_t max_msg_size = 1 << 20;
+  size_t min_msg_size = 1;
+  size_t max_msg_size = 1 << 20;
   unsigned wg_size = 64;
   unsigned thread_access = 64;
   unsigned coal_coef = 64;
@@ -74,7 +74,7 @@ class TesterArguments {
   int loop = 10;
   int skip = 10;
   int loop_large = 10;
-  uint64_t large_message_size = 32768;
+  size_t large_message_size = 32768;
 };
 
 #endif

tests/functional_tests/wavefront_primitives.cpp

@@ -36,7 +36,7 @@ using namespace rocshmem;
 __global__ void WaveFrontPrimitiveTest(int loop, int skip,
                                        long long int *start_time,
                                        long long int *end_time, char *source,
-                                       char *dest, int size, TestType type,
+                                       char *dest, size_t size, TestType type,
                                        ShmemContextType ctx_type,
                                        int wf_size) {
   __shared__ rocshmem_ctx_t ctx;
@@ -49,7 +49,7 @@ __global__ void WaveFrontPrimitiveTest(int loop, int skip,
   int wf_id = get_flat_block_id() / wf_size;
   int wg_offset = wg_id * ((get_flat_block_size() - 1 ) / wf_size + 1);
   int idx = wf_id + wg_offset;
-  int offset = size * idx;
+  size_t offset = size * idx;
   source += offset;
   dest += offset;
 
@@ -120,13 +120,13 @@ WaveFrontPrimitiveTester::~WaveFrontPrimitiveTester() {
   rocshmem_free(dest);
 }
 
-void WaveFrontPrimitiveTester::resetBuffers(uint64_t size) {
+void WaveFrontPrimitiveTester::resetBuffers(size_t size) {
   size_t buff_size = size * args.num_wgs * num_warps;
   memset(dest, '1', buff_size);
 }
 
 void WaveFrontPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize,
-                                           int loop, uint64_t size) {
+                                           int loop, size_t size) {
   size_t shared_bytes = 0;
 
   hipLaunchKernelGGL(WaveFrontPrimitiveTest, gridSize, blockSize, shared_bytes,
@@ -138,7 +138,7 @@ void WaveFrontPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize,
   num_timed_msgs = loop * gridSize.x * num_warps;
 }
 
-void WaveFrontPrimitiveTester::verifyResults(uint64_t size) {
+void WaveFrontPrimitiveTester::verifyResults(size_t size) {
   int check_id = (_type == WAVEGetTestType || _type == WAVEGetNBITestType)
                      ? 0
                      : 1;

tests/functional_tests/wavefront_primitives.hpp

@@ -36,12 +36,12 @@ class WaveFrontPrimitiveTester : public Tester {
   virtual ~WaveFrontPrimitiveTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   char *source = nullptr;
   char *dest = nullptr;

tests/functional_tests/workgroup_primitives.cpp

@@ -36,7 +36,7 @@ using namespace rocshmem;
 __global__ void WorkGroupPrimitiveTest(int loop, int skip,
                                       long long int *start_time,
                                       long long int *end_time, char *source,
-                                      char *dest, int size, TestType type,
+                                      char *dest, size_t size, TestType type,
                                       ShmemContextType ctx_type) {
   __shared__ rocshmem_ctx_t ctx;
   int wg_id = get_flat_grid_id();
@@ -44,7 +44,7 @@ __global__ void WorkGroupPrimitiveTest(int loop, int skip,
   rocshmem_wg_ctx_create(ctx_type, &ctx);
 
   // Calculate start index for each work group
-  uint64_t offset = size * wg_id;
+  size_t offset = size * wg_id;
   source += offset;
   dest += offset;
 
@@ -116,13 +116,13 @@ WorkGroupPrimitiveTester::~WorkGroupPrimitiveTester() {
   rocshmem_free(dest);
 }
 
-void WorkGroupPrimitiveTester::resetBuffers(uint64_t size) {
+void WorkGroupPrimitiveTester::resetBuffers(size_t size) {
   size_t buff_size = size * args.num_wgs;
   memset(dest, '1', buff_size);
 }
 
 void WorkGroupPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize,
-                                           int loop, uint64_t size) {
+                                           int loop, size_t size) {
   size_t shared_bytes = 0;
 
   hipLaunchKernelGGL(WorkGroupPrimitiveTest, gridSize, blockSize, shared_bytes,
@@ -133,7 +133,7 @@ void WorkGroupPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize,
   num_timed_msgs = loop * gridSize.x;
 }
 
-void WorkGroupPrimitiveTester::verifyResults(uint64_t size) {
+void WorkGroupPrimitiveTester::verifyResults(size_t size) {
   int check_id = (_type == WGGetTestType || _type == WGGetNBITestType)
                      ? 0
                      : 1;

tests/functional_tests/workgroup_primitives.hpp

@@ -36,12 +36,12 @@ class WorkGroupPrimitiveTester : public Tester {
   virtual ~WorkGroupPrimitiveTester();
 
  protected:
-  virtual void resetBuffers(uint64_t size) override;
+  virtual void resetBuffers(size_t size) override;
 
   virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
-                            uint64_t size) override;
+                            size_t size) override;
 
-  virtual void verifyResults(uint64_t size) override;
+  virtual void verifyResults(size_t size) override;
 
   char *source = nullptr;
   char *dest = nullptr;