diff --git a/projects/rocshmem/CHANGELOG.md b/projects/rocshmem/CHANGELOG.md index bd7b75fc53..a268d061d7 100644 --- a/projects/rocshmem/CHANGELOG.md +++ b/projects/rocshmem/CHANGELOG.md @@ -1,5 +1,14 @@ # Changelog for rocSHMEM +## rocSHMEM 3.x.x for ROCm 7.x.x + +### Changed + +* The following APIs have been deprecated: + * `rocshmem_wg_init` + * `rocshmem_wg_finalize` + * `rocshmem_wg_init_thread` + ## rocSHMEM 3.0.0 for ROCm 7.0.0 ### Added diff --git a/projects/rocshmem/docs/api/init.rst b/projects/rocshmem/docs/api/init.rst index e381ca5663..f605b729ad 100644 --- a/projects/rocshmem/docs/api/init.rst +++ b/projects/rocshmem/docs/api/init.rst @@ -23,12 +23,16 @@ you must select the device that this PE is associated to by calling `hipSetDevice `_. -.. cpp:function:: __device__ void rocshmem_wg_init(void) +.. WARNING:: + Routine `rocshmem_wg_init` has been deprecated. + +.. cpp:function:: [[deprecated]] __device__ void rocshmem_wg_init(void) :Parameters: None. :returns: None. **Description:** +This routine has been deprecated, please do not use. This routine initializes device-side rocSHMEM resources. It must be called before any threads in this work-group invoke other rocSHMEM functions. It must be called collectively by all threads in the work-group. @@ -43,12 +47,16 @@ ROCSHMEM_FINALIZE **Description:** This routine finalizes the rocSHMEM library. -.. cpp:function:: __device__ void rocshmem_wg_finalize(void) +.. WARNING:: + Routine `rocshmem_wg_finalize` has been deprecated. + +.. cpp:function:: [[deprecated]] __device__ void rocshmem_wg_finalize(void) :Parameters: None. :returns: None. **Description:** +This routine has been deprecated, please do not use. This routine finalizes device-side rocSHMEM resources. It must be called before work-group completion if the work-group also called ``rocshmem_wg_init``. It must be called collectively by all threads in the work-group. @@ -65,7 +73,7 @@ ROCSHMEM_INIT_ATTR **Description:** This routine initializes the rocSHMEM runtime and underlying transport layer using the provided mode and attributes. -The parameter ``flags`` can be either +The parameter ``flags`` can be either ``ROCSHMEM_INIT_WITH_UNIQUEID`` or ``ROCSHMEM_INIT_WITH_MPI_COMM``. ROCSHMEM_GET_UNIQUEID diff --git a/projects/rocshmem/examples/rocshmem_allreduce_test.cc b/projects/rocshmem/examples/rocshmem_allreduce_test.cc index 50f4483784..943087a511 100644 --- a/projects/rocshmem/examples/rocshmem_allreduce_test.cc +++ b/projects/rocshmem/examples/rocshmem_allreduce_test.cc @@ -65,7 +65,6 @@ __global__ void allreduce_test(int *source, int *dest, size_t nelem, __shared__ rocshmem_ctx_t ctx; int64_t ctx_type = 0; - rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); int num_pes = rocshmem_ctx_n_pes(ctx); @@ -75,7 +74,6 @@ __global__ void allreduce_test(int *source, int *dest, size_t nelem, __syncthreads(); rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } static void init_sendbuf (int *source, int nelem, int my_pe) diff --git a/projects/rocshmem/examples/rocshmem_alltoall_test.cc b/projects/rocshmem/examples/rocshmem_alltoall_test.cc index 90d7cc9437..3719e42372 100644 --- a/projects/rocshmem/examples/rocshmem_alltoall_test.cc +++ b/projects/rocshmem/examples/rocshmem_alltoall_test.cc @@ -65,7 +65,6 @@ __global__ void alltoall_test(int *source, int *dest, size_t nelem, __shared__ rocshmem_ctx_t ctx; int64_t ctx_type = 0; - rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); int num_pes = rocshmem_ctx_n_pes(ctx); @@ -75,7 +74,6 @@ __global__ void alltoall_test(int *source, int *dest, size_t nelem, __syncthreads(); rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } static void init_sendbuf (int *source, int nelem, int my_pe, int npes) diff --git a/projects/rocshmem/examples/rocshmem_broadcast_test.cc b/projects/rocshmem/examples/rocshmem_broadcast_test.cc index 7fa895bed5..f10324e94b 100644 --- a/projects/rocshmem/examples/rocshmem_broadcast_test.cc +++ b/projects/rocshmem/examples/rocshmem_broadcast_test.cc @@ -65,7 +65,6 @@ __global__ void broadcast_test(int *source, int *dest, size_t nelem, __shared__ rocshmem_ctx_t ctx; int64_t ctx_type = 0; - rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); int num_pes = rocshmem_ctx_n_pes(ctx); @@ -75,7 +74,6 @@ __global__ void broadcast_test(int *source, int *dest, size_t nelem, __syncthreads(); rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } static void init_sendbuf(int *source, int nelem, int my_pe) diff --git a/projects/rocshmem/examples/rocshmem_getmem_test.cc b/projects/rocshmem/examples/rocshmem_getmem_test.cc index 3c77dd43ab..ea2e9f625d 100644 --- a/projects/rocshmem/examples/rocshmem_getmem_test.cc +++ b/projects/rocshmem/examples/rocshmem_getmem_test.cc @@ -62,7 +62,6 @@ using namespace rocshmem; __global__ void simple_getmem_test(int *src, int *dst, size_t nelem) { - rocshmem_wg_init(); int threadId = blockIdx.x * blockDim.x + threadIdx.x; if (threadId == 0) { @@ -73,7 +72,6 @@ __global__ void simple_getmem_test(int *src, int *dst, size_t nelem) } __syncthreads(); - rocshmem_wg_finalize(); } #define MAX_ELEM 256 diff --git a/projects/rocshmem/examples/rocshmem_put_signal_test.cc b/projects/rocshmem/examples/rocshmem_put_signal_test.cc index 27086f4c07..a23e164657 100644 --- a/projects/rocshmem/examples/rocshmem_put_signal_test.cc +++ b/projects/rocshmem/examples/rocshmem_put_signal_test.cc @@ -63,7 +63,6 @@ using namespace rocshmem; __global__ void simple_put_signal_test(uint64_t *data, uint64_t *message, size_t nelem, uint64_t *sig_addr, int my_pe, int dst_pe) { - rocshmem_wg_init(); int threadId = blockIdx.x * blockDim.x + threadIdx.x; @@ -78,7 +77,6 @@ __global__ void simple_put_signal_test(uint64_t *data, uint64_t *message, size_t } __syncthreads(); - rocshmem_wg_finalize(); } #define MAX_ELEM 256 diff --git a/projects/rocshmem/include/rocshmem/rocshmem.hpp b/projects/rocshmem/include/rocshmem/rocshmem.hpp index 4e30438388..c42c7fbe60 100644 --- a/projects/rocshmem/include/rocshmem/rocshmem.hpp +++ b/projects/rocshmem/include/rocshmem/rocshmem.hpp @@ -374,7 +374,7 @@ __host__ void rocshmem_global_exit(int status); * * @return void. */ -__device__ void rocshmem_wg_init(); +[[deprecated]] __device__ void rocshmem_wg_init(); /** * @brief Finalizes device-side rocSHMEM resources. Must be called before @@ -384,7 +384,7 @@ __device__ void rocshmem_wg_init(); * * @return void. */ -__device__ void rocshmem_wg_finalize(); +[[deprecated]] __device__ void rocshmem_wg_finalize(); /** * @brief Initializes device-side rocSHMEM resources. Must be called before @@ -400,7 +400,7 @@ __device__ void rocshmem_wg_finalize(); * * @return void. */ -__device__ void rocshmem_wg_init_thread(int requested, int *provided); +[[deprecated]] __device__ void rocshmem_wg_init_thread(int requested, int *provided); /** * @brief Query the thread mode used by the runtime. diff --git a/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp b/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp index 25c8c6dfe5..4912fc3c72 100644 --- a/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp @@ -254,7 +254,6 @@ void AMOBitwiseTester::verifyResults(size_t size) { int global_id = get_flat_id(); \ int n_threads = get_flat_grid_size(); \ int n_wgs = get_grid_num_blocks(); \ - rocshmem_wg_init(); \ rocshmem_wg_ctx_create(ctx_type, &ctx); \ for (int i = 0; i < loop + skip; i++) { \ T *ptr = compute_target_ptr(dest, addr_mode, wg_id, i, n_wgs); \ @@ -293,7 +292,6 @@ void AMOBitwiseTester::verifyResults(size_t size) { end_time[wg_id] = wall_clock64(); \ __syncthreads(); \ rocshmem_wg_ctx_destroy(&ctx); \ - rocshmem_wg_finalize(); \ } \ template class AMOBitwiseTester; diff --git a/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp b/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp index b4bc76f754..630069ce28 100644 --- a/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp @@ -229,7 +229,6 @@ void AMOExtendedTester::verifyResults(size_t /*size*/) { int t_id = get_flat_block_id(); \ int n_threads = get_flat_grid_size(); \ int n_wgs = get_grid_num_blocks(); \ - rocshmem_wg_init(); \ rocshmem_wg_ctx_create(ctx_type, &ctx); \ for (int i = 0; i < loop + skip; i++) { \ T *ptr = compute_target_ptr(dest, addr_mode, wg_id, i, n_wgs); \ @@ -256,7 +255,6 @@ void AMOExtendedTester::verifyResults(size_t /*size*/) { end_time[wg_id] = wall_clock64(); \ __syncthreads(); \ rocshmem_wg_ctx_destroy(&ctx); \ - rocshmem_wg_finalize(); \ } \ template class AMOExtendedTester; diff --git a/projects/rocshmem/tests/functional_tests/amo_standard_tester.cpp b/projects/rocshmem/tests/functional_tests/amo_standard_tester.cpp index 7b81d36f1a..0e30b1f5c7 100644 --- a/projects/rocshmem/tests/functional_tests/amo_standard_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/amo_standard_tester.cpp @@ -260,7 +260,6 @@ void AMOStandardTester::verifyResults(size_t size) { int t_id = get_flat_block_id(); \ int n_threads = get_flat_grid_size(); \ int n_wgs = get_grid_num_blocks(); \ - rocshmem_wg_init(); \ rocshmem_wg_ctx_create(ctx_type, &ctx); \ for (int i = 0; i < loop + skip; i++) { \ T *ptr = compute_target_ptr(dest, addr_mode, wg_id, i, n_wgs); \ @@ -294,7 +293,6 @@ void AMOStandardTester::verifyResults(size_t size) { end_time[wg_id] = wall_clock64(); \ __syncthreads(); \ rocshmem_wg_ctx_destroy(&ctx); \ - rocshmem_wg_finalize(); \ } \ template class AMOStandardTester; diff --git a/projects/rocshmem/tests/functional_tests/barrier_all_tester.cpp b/projects/rocshmem/tests/functional_tests/barrier_all_tester.cpp index 260d73319a..62798099a2 100644 --- a/projects/rocshmem/tests/functional_tests/barrier_all_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/barrier_all_tester.cpp @@ -39,7 +39,6 @@ __global__ void BarrierAllTest(int loop, int skip, long long int *start_time, int wg_id = get_flat_grid_id(); int wf_id = t_id / wf_size; - rocshmem_wg_init(); for (int i = 0; i < loop + skip; i++) { if (hipThreadIdx_x == 0 && i == skip) { @@ -84,7 +83,6 @@ __global__ void BarrierAllTest(int loop, int skip, long long int *start_time, end_time[wg_id] = wall_clock64(); } - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/default_ctx_primitive_tester.cpp b/projects/rocshmem/tests/functional_tests/default_ctx_primitive_tester.cpp index e64ae844ca..f8b3ce4db5 100644 --- a/projects/rocshmem/tests/functional_tests/default_ctx_primitive_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/default_ctx_primitive_tester.cpp @@ -39,7 +39,6 @@ int wg_id = get_flat_grid_id(); int t_id = get_flat_block_id(); int wf_id = t_id / wf_size; - rocshmem_wg_init(); /** * Shared array to capture the start time for each wavefront @@ -121,7 +120,6 @@ start_time[wg_id] = wf_start_time[0]; } - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/empty_tester.cpp b/projects/rocshmem/tests/functional_tests/empty_tester.cpp index f244d27712..83048378ac 100644 --- a/projects/rocshmem/tests/functional_tests/empty_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/empty_tester.cpp @@ -35,11 +35,9 @@ __global__ void EmptyTest(int loop, int skip, long long int *start_time, long long int *end_time, int size, TestType type, ShmemContextType ctx_type) { __shared__ rocshmem_ctx_t ctx; - rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/ping_all_tester.cpp b/projects/rocshmem/tests/functional_tests/ping_all_tester.cpp index 0bdcc0b840..ed7abd1d22 100644 --- a/projects/rocshmem/tests/functional_tests/ping_all_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/ping_all_tester.cpp @@ -37,7 +37,6 @@ __global__ void PingAllTest(int loop, int skip, long long int *start_time, __shared__ rocshmem_ctx_t ctx; int wg_id = get_flat_grid_id(); - rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); int pe = rocshmem_ctx_my_pe(ctx); @@ -64,7 +63,6 @@ __global__ void PingAllTest(int loop, int skip, long long int *start_time, rocshmem_ctx_quiet(ctx); } rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/ping_pong_tester.cpp b/projects/rocshmem/tests/functional_tests/ping_pong_tester.cpp index c3e9ead963..2b7b319abc 100644 --- a/projects/rocshmem/tests/functional_tests/ping_pong_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/ping_pong_tester.cpp @@ -37,7 +37,6 @@ __global__ void PingPongTest(int loop, int skip, long long int *start_time, __shared__ rocshmem_ctx_t ctx; int wg_id = get_flat_grid_id(); - rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); int pe = rocshmem_ctx_my_pe(ctx); @@ -65,7 +64,6 @@ __global__ void PingPongTest(int loop, int skip, long long int *start_time, } rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/primitive_mr_tester.cpp b/projects/rocshmem/tests/functional_tests/primitive_mr_tester.cpp index 22483f73b7..d5df5afc40 100644 --- a/projects/rocshmem/tests/functional_tests/primitive_mr_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/primitive_mr_tester.cpp @@ -37,7 +37,6 @@ __global__ void PrimitiveMRTest(int loop, long long int *start_time, ShmemContextType ctx_type) { __shared__ rocshmem_ctx_t ctx; int wg_id = get_flat_grid_id(); - rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); if (hipThreadIdx_x == 0) { @@ -57,7 +56,6 @@ __global__ void PrimitiveMRTest(int loop, long long int *start_time, __syncthreads(); rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/primitive_tester.cpp b/projects/rocshmem/tests/functional_tests/primitive_tester.cpp index 83b3953f42..6cef5e39b4 100644 --- a/projects/rocshmem/tests/functional_tests/primitive_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/primitive_tester.cpp @@ -39,7 +39,6 @@ __global__ void PrimitiveTest(int loop, int skip, long long int *start_time, int wg_id = get_flat_grid_id(); int t_id = get_flat_block_id(); int wf_id = t_id / wf_size; - rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); /** @@ -121,7 +120,6 @@ __global__ void PrimitiveTest(int loop, int skip, long long int *start_time, } rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/random_access_tester.cpp b/projects/rocshmem/tests/functional_tests/random_access_tester.cpp index b195259e76..20274b3708 100644 --- a/projects/rocshmem/tests/functional_tests/random_access_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/random_access_tester.cpp @@ -62,7 +62,6 @@ __global__ void RandomAccessTest(int loop, int skip, long long int *start_time, uint32_t *PE_bins, ShmemContextType ctx_type) { __shared__ rocshmem_ctx_t ctx; int wg_id = get_flat_grid_id(); - rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); int pe = rocshmem_ctx_my_pe(ctx); @@ -97,7 +96,6 @@ __global__ void RandomAccessTest(int loop, int skip, long long int *start_time, end_time[wg_id] = wall_clock64(); } rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/shmem_ptr_tester.cpp b/projects/rocshmem/tests/functional_tests/shmem_ptr_tester.cpp index 58dbaeab3c..158b6d0878 100644 --- a/projects/rocshmem/tests/functional_tests/shmem_ptr_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/shmem_ptr_tester.cpp @@ -39,7 +39,6 @@ __global__ void ShmemPtrTest(int loop, int skip, long long int *start_time, int t_id = get_flat_block_id(); int wf_id = t_id / wf_size; - rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); /** @@ -111,7 +110,6 @@ __global__ void ShmemPtrTest(int loop, int skip, long long int *start_time, } rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/signaling_operations_tester.cpp b/projects/rocshmem/tests/functional_tests/signaling_operations_tester.cpp index ff37945cf0..2aa73ec32e 100644 --- a/projects/rocshmem/tests/functional_tests/signaling_operations_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/signaling_operations_tester.cpp @@ -38,7 +38,6 @@ __global__ void PutmemSignalTest(int loop, int skip, long long int *start_time, int sig_op) { __shared__ rocshmem_ctx_t ctx; int wg_id = get_flat_grid_id(); - rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); uint64_t signal = 1; @@ -88,13 +87,11 @@ __global__ void PutmemSignalTest(int loop, int skip, long long int *start_time, } rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } __global__ void SignalFetchTest(int loop, int skip, long long int *start_time, long long int *end_time, uint64_t *sig_addr, uint64_t *fetched_value, TestType type) { - rocshmem_wg_init(); int wg_id = get_flat_grid_id(); @@ -125,7 +122,6 @@ __global__ void SignalFetchTest(int loop, int skip, long long int *start_time, end_time[wg_id] = wall_clock64(); } - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/sync_all_tester.cpp b/projects/rocshmem/tests/functional_tests/sync_all_tester.cpp index df26626041..df797d51ea 100644 --- a/projects/rocshmem/tests/functional_tests/sync_all_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/sync_all_tester.cpp @@ -39,7 +39,6 @@ __global__ void SyncAllTest(int loop, int skip, long long int *start_time, int wg_id = get_flat_grid_id(); int wf_id = t_id / wf_size; - rocshmem_wg_init(); for (int i = 0; i < loop + skip; i++) { if (hipThreadIdx_x == 0 && i == skip) { @@ -84,7 +83,6 @@ __global__ void SyncAllTest(int loop, int skip, long long int *start_time, end_time[wg_id] = wall_clock64(); } - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/sync_tester.cpp b/projects/rocshmem/tests/functional_tests/sync_tester.cpp index 9bd9bd6f0a..1e7c583523 100644 --- a/projects/rocshmem/tests/functional_tests/sync_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/sync_tester.cpp @@ -36,7 +36,6 @@ __global__ void SyncTest(int loop, int skip, long long int *start_time, int wg_id = get_flat_grid_id(); int wf_id = t_id / wf_size; - rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); for (int i = 0; i < loop + skip; i++) { @@ -69,7 +68,6 @@ __global__ void SyncTest(int loop, int skip, long long int *start_time, } rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/team_alltoall_tester.cpp b/projects/rocshmem/tests/functional_tests/team_alltoall_tester.cpp index 6c11859a81..3a8c964701 100644 --- a/projects/rocshmem/tests/functional_tests/team_alltoall_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/team_alltoall_tester.cpp @@ -64,7 +64,6 @@ __global__ void TeamAlltoallTest(int loop, int skip, long long int *start_time, __shared__ rocshmem_ctx_t ctx; int wg_id = get_flat_grid_id(); - rocshmem_wg_init(); rocshmem_wg_team_create_ctx(teams[wg_id], ctx_type, &ctx); int n_pes = rocshmem_ctx_n_pes(ctx); @@ -91,7 +90,6 @@ __global__ void TeamAlltoallTest(int loop, int skip, long long int *start_time, } rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/team_barrier_tester.cpp b/projects/rocshmem/tests/functional_tests/team_barrier_tester.cpp index bd0430537d..dc3ccf61c1 100644 --- a/projects/rocshmem/tests/functional_tests/team_barrier_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/team_barrier_tester.cpp @@ -36,7 +36,6 @@ __global__ void TeamBarrierTest(int loop, int skip, long long int *start_time, int wg_id = get_flat_grid_id(); int wf_id = t_id / wf_size; - rocshmem_wg_init(); rocshmem_wg_team_create_ctx(teams[wg_id], ctx_type, &ctx); for (int i = 0; i < loop + skip; i++) { @@ -69,7 +68,6 @@ __global__ void TeamBarrierTest(int loop, int skip, long long int *start_time, } rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/team_broadcast_tester.cpp b/projects/rocshmem/tests/functional_tests/team_broadcast_tester.cpp index 393916a814..b747508608 100644 --- a/projects/rocshmem/tests/functional_tests/team_broadcast_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/team_broadcast_tester.cpp @@ -69,7 +69,6 @@ __global__ void TeamBroadcastTest(int loop, int skip, long long int *start_time, __shared__ rocshmem_ctx_t ctx; int wg_id = get_flat_grid_id(); - rocshmem_wg_init(); rocshmem_wg_team_create_ctx(teams[wg_id], ctx_type, &ctx); int n_pes = rocshmem_ctx_n_pes(ctx); @@ -97,7 +96,6 @@ __global__ void TeamBroadcastTest(int loop, int skip, long long int *start_time, } rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/team_ctx_infra_tester.cpp b/projects/rocshmem/tests/functional_tests/team_ctx_infra_tester.cpp index aeeb479ae9..8421db437a 100644 --- a/projects/rocshmem/tests/functional_tests/team_ctx_infra_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/team_ctx_infra_tester.cpp @@ -44,7 +44,6 @@ rocshmem_team_t team_world_dup[NUM_TEAMS]; int expected_pe, int expected_n_pes) { __shared__ rocshmem_ctx_t ctx; - rocshmem_wg_init(); rocshmem_wg_team_create_ctx(team, ctx_type, &ctx); int num_pes = rocshmem_ctx_n_pes(ctx); @@ -64,7 +63,6 @@ rocshmem_team_t team_world_dup[NUM_TEAMS]; rocshmem_ctx_quiet(ctx); rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } __global__ void TeamCtxInfraTest(ShmemContextType ctx_type, @@ -72,7 +70,6 @@ rocshmem_team_t team_world_dup[NUM_TEAMS]; __shared__ rocshmem_ctx_t ctx1, ctx2, ctx3; __shared__ rocshmem_ctx_t ctx[NUM_TEAMS]; - rocshmem_wg_init(); /** * Test 1: Assert team infos of different ctxs @@ -131,7 +128,6 @@ rocshmem_team_t team_world_dup[NUM_TEAMS]; rocshmem_wg_ctx_destroy(&ctx[team_i]); } - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/team_ctx_primitive_tester.cpp b/projects/rocshmem/tests/functional_tests/team_ctx_primitive_tester.cpp index 8e5e10ea97..673a536d58 100644 --- a/projects/rocshmem/tests/functional_tests/team_ctx_primitive_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/team_ctx_primitive_tester.cpp @@ -43,7 +43,6 @@ __global__ void TeamCtxPrimitiveTest(int loop, int skip, long long int *start_ti int t_id = get_flat_block_id(); int wf_id = t_id / wf_size; - rocshmem_wg_init(); rocshmem_wg_team_create_ctx(team, ctx_type, &ctx); /** @@ -114,7 +113,6 @@ __global__ void TeamCtxPrimitiveTest(int loop, int skip, long long int *start_ti } rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/team_fcollect_tester.cpp b/projects/rocshmem/tests/functional_tests/team_fcollect_tester.cpp index 244556e25d..73cbc28f03 100644 --- a/projects/rocshmem/tests/functional_tests/team_fcollect_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/team_fcollect_tester.cpp @@ -64,7 +64,6 @@ __global__ void TeamFcollectTest(int loop, int skip, long long int *start_time, __shared__ rocshmem_ctx_t ctx; int wg_id = get_flat_grid_id(); - rocshmem_wg_init(); rocshmem_wg_team_create_ctx(teams[wg_id], ctx_type, &ctx); int n_pes = rocshmem_ctx_n_pes(ctx); @@ -90,7 +89,6 @@ __global__ void TeamFcollectTest(int loop, int skip, long long int *start_time, } rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/team_reduction_tester.cpp b/projects/rocshmem/tests/functional_tests/team_reduction_tester.cpp index a826678fd2..ec19bee00a 100644 --- a/projects/rocshmem/tests/functional_tests/team_reduction_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/team_reduction_tester.cpp @@ -83,7 +83,6 @@ __global__ void TeamReductionTest(int loop, int skip, long long int *start_time, __shared__ rocshmem_ctx_t ctx; int wg_id = get_flat_grid_id(); - rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); int n_pes = rocshmem_ctx_n_pes(ctx); @@ -104,7 +103,6 @@ __global__ void TeamReductionTest(int loop, int skip, long long int *start_time, } rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/wavefront_primitives.cpp b/projects/rocshmem/tests/functional_tests/wavefront_primitives.cpp index be18c1e822..14bbe572b0 100644 --- a/projects/rocshmem/tests/functional_tests/wavefront_primitives.cpp +++ b/projects/rocshmem/tests/functional_tests/wavefront_primitives.cpp @@ -42,7 +42,6 @@ __global__ void WaveFrontPrimitiveTest(int loop, int skip, __shared__ rocshmem_ctx_t ctx; int wg_id = get_flat_grid_id(); - rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); // Calculate start index for each wavefront @@ -86,7 +85,6 @@ __global__ void WaveFrontPrimitiveTest(int loop, int skip, } rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/workgroup_primitives.cpp b/projects/rocshmem/tests/functional_tests/workgroup_primitives.cpp index f4e14b6d4d..d0f4a0fd3d 100644 --- a/projects/rocshmem/tests/functional_tests/workgroup_primitives.cpp +++ b/projects/rocshmem/tests/functional_tests/workgroup_primitives.cpp @@ -40,7 +40,6 @@ __global__ void WorkGroupPrimitiveTest(int loop, int skip, ShmemContextType ctx_type) { __shared__ rocshmem_ctx_t ctx; int wg_id = get_flat_grid_id(); - rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); // Calculate start index for each work group @@ -82,7 +81,6 @@ __global__ void WorkGroupPrimitiveTest(int loop, int skip, } rocshmem_wg_ctx_destroy(&ctx); - rocshmem_wg_finalize(); } /******************************************************************************