Add host API for *_on_stream operations (#340)
* Add functional test for barrier_all_on_stream
* Add rocshmem_barrier_all_on_stream support for GDA and RO backends
Implements rocshmem_barrier_all_on_stream operation for
GPU Direct Access and Reverse Offload backends.
Previously, rocshmem_barrier_all_on_stream was only supported for IPC backend.
* Add functional test for rocshmem_broadcastmem_on_stream
* Add host-side rocshmem_broadcastmem_on_stream API
Implement stream-based broadcast collective operation
- Add rocshmem_broadcastmem_on_stream host API and kernel implementation
- Add functional test TeamBroadcastmemOnStreamTester with multi-stream
support and correctness verification
- Use per-workgroup contexts to avoid contention across parallel streams
API:
rocshmem_broadcastmem_on_stream(team, dest, source, nelems, pe_root, stream)
* Add functional test for rocshmem_getmem_on_stream
* Add host-side rocshmem_getmem_on_stream API
Implement stream-based point-to-point RMA get operation
- Add rocshmem_getmem_on_stream host API and kernel implementation
- Support for asynchronous getmem operations on HIP streams
- Add backend support for GDA, RO, and IPC contexts
- Use work-group collective getmem for efficient memory transfer
API:
rocshmem_getmem_on_stream(dest, source, nelems, pe, stream)
(AI Assist)
* Add host-side rocshmem_putmem_on_stream API
- Add rocshmem_putmem_on_stream for asynchronous remote writes
- Support for concurrent RMA operations on HIP streams
- Add backend support for GDA, RO, and IPC contexts
- Use work-group device collective operation
API:
rocshmem_putmem_on_stream(dest, source, bytes, pe, stream)
(AI Assist)
* Add functional test for rocshmem_putmem_on_stream
* Add host-side rocshmem_putmem_signal_on_stream API
Enables asynchronous putmem operations with signaling on HIP streams.
The implementation includes:
- Kernel wrapper rocshmem_putmem_signal_kernel
- Host interface putmem_signal_on_stream method
- Context layer support across all backends (IPC, GDA, RO)
- Public API
Function signature:
void rocshmem_putmem_signal_on_stream(void *dest, const void *source,
size_t bytes, uint64_t *sig_addr,
uint64_t signal, int sig_op,
int pe, hipStream_t stream);
* Add functional test for rocshmem_putmem_signal_on_stream
* Add host-side rocshmem_signal_wait_until_on_stream API
Enables asynchronous signal wait operations on HIP streams.
The implementation includes:
- Kernel wrapper rocshmem_signal_wait_until_kernel
- Host interface signal_wait_until_on_stream method
- Context layer support across all backends (IPC, GDA, RO)
- Native uint64_t support in wait_until API (generated from P2P_SYNC.py)
Function signature:
void rocshmem_signal_wait_until_on_stream(uint64_t *sig_addr, int cmp,
uint64_t cmp_value,
hipStream_t stream);
(AI Assist)
* Add functional test for rocshmem_signal_wait_until_on_stream
* Add documentation for stream API functions
This commit adds API documentation for the following host-side
stream functions:
- rocshmem_barrier_all_on_stream (collective routines)
- rocshmem_broadcastmem_on_stream (collective routines)
- rocshmem_getmem_on_stream (RMA operations)
- rocshmem_putmem_on_stream (RMA operations)
- rocshmem_putmem_signal_on_stream (signaling operations)
- rocshmem_signal_wait_until_on_stream (point-to-point sync)
The documentation includes function signatures, parameter descriptions,
and detailed explanations of asynchronous behavior and stream handling.
(AI Assist)
* Rename "bytes" -> "nelems"
* Add "_TEST_" to the variables used in tests
* Remove incorrect hipStreamDefault usage
hipStreamDefault is not a default stream. This is a flag.
If stream == nullptr, then just pass it to kernel. It will launch the kernel on the default stream
This commit is contained in:
Anatolii Rozanov
GitHub
@@ -0,0 +1,151 @@
|
||||
/******************************************************************************
|
||||
* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* SPDX-License-Identifier: MIT
|
||||
*
|
||||
* 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 "barrier_all_on_stream_tester.hpp"
|
||||
|
||||
#include <rocshmem/rocshmem.hpp>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <cstring>
|
||||
#include <cassert>
|
||||
#include <vector>
|
||||
|
||||
/******************************************************************************
|
||||
* HOST TESTER CLASS METHODS
|
||||
*****************************************************************************/
|
||||
BarrierAllOnStreamTester::BarrierAllOnStreamTester(TesterArguments args)
|
||||
: Tester(args) {
|
||||
my_pe = rocshmem_my_pe();
|
||||
n_pes = rocshmem_n_pes();
|
||||
|
||||
char *value{nullptr};
|
||||
if ((value = getenv("ROCSHMEM_TEST_NUM_STREAMS"))) {
|
||||
num_streams = atoi(value);
|
||||
} else {
|
||||
// Default to 1 stream
|
||||
num_streams = 1;
|
||||
}
|
||||
|
||||
// Check if we should test with nullptr (default stream)
|
||||
use_default_stream = false;
|
||||
if ((value = getenv("ROCSHMEM_TEST_USE_DEFAULT_STREAM"))) {
|
||||
use_default_stream = (atoi(value) != 0);
|
||||
if (use_default_stream) {
|
||||
num_streams = 1; // Only test with one nullptr stream
|
||||
}
|
||||
}
|
||||
|
||||
streams.resize(num_streams);
|
||||
start_events_timed.resize(num_streams);
|
||||
stop_events_timed.resize(num_streams);
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
if (use_default_stream) {
|
||||
streams[i] = nullptr; // Use default stream (0)
|
||||
} else {
|
||||
CHECK_HIP(hipStreamCreate(&streams[i]));
|
||||
}
|
||||
CHECK_HIP(hipEventCreate(&start_events_timed[i]));
|
||||
CHECK_HIP(hipEventCreate(&stop_events_timed[i]));
|
||||
}
|
||||
}
|
||||
|
||||
BarrierAllOnStreamTester::~BarrierAllOnStreamTester() {
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipEventDestroy(stop_events_timed[i]));
|
||||
CHECK_HIP(hipEventDestroy(start_events_timed[i]));
|
||||
// Don't destroy default stream (nullptr)
|
||||
if (!use_default_stream) {
|
||||
CHECK_HIP(hipStreamDestroy(streams[i]));
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
void BarrierAllOnStreamTester::preLaunchKernel() {
|
||||
// No specific setup needed for barrier
|
||||
}
|
||||
|
||||
void BarrierAllOnStreamTester::postLaunchKernel() {
|
||||
// Synchronize all streams to ensure events are recorded
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipStreamSynchronize(streams[i]));
|
||||
}
|
||||
|
||||
// Get elapsed time for each stream from HIP events
|
||||
for (int stream_id = 0; stream_id < num_streams && stream_id < num_timers;
|
||||
stream_id++) {
|
||||
float elapsed_time_ms = 0.0f;
|
||||
CHECK_HIP(hipEventElapsedTime(&elapsed_time_ms,
|
||||
start_events_timed[stream_id],
|
||||
stop_events_timed[stream_id]));
|
||||
|
||||
// Convert milliseconds to GPU cycles
|
||||
// wall_clk_rate is in kHz, so: cycles = ms * wall_clk_rate
|
||||
long long int elapsed_cycles =
|
||||
static_cast<long long int>(elapsed_time_ms *
|
||||
static_cast<float>(wall_clk_rate));
|
||||
|
||||
start_time[stream_id] = 0;
|
||||
end_time[stream_id] = elapsed_cycles;
|
||||
}
|
||||
|
||||
// Fill remaining timers with zero if num_timers > num_streams
|
||||
for (int i = num_streams; i < num_timers; i++) {
|
||||
start_time[i] = 0;
|
||||
end_time[i] = 0;
|
||||
}
|
||||
}
|
||||
|
||||
void BarrierAllOnStreamTester::resetBuffers(size_t size) {}
|
||||
|
||||
void BarrierAllOnStreamTester::launchKernel(dim3 gridSize, dim3 blockSize,
|
||||
int loop, size_t size) {
|
||||
// Execute warmup iterations (skip)
|
||||
for (int i = 0; i < args.skip; i++) {
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
rocshmem_barrier_all_on_stream(streams[stream_id]);
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < loop; i++) {
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
// Record start event for this stream on first iteration
|
||||
if (i == 0) {
|
||||
CHECK_HIP(hipEventRecord(start_events_timed[stream_id],
|
||||
streams[stream_id]));
|
||||
}
|
||||
|
||||
rocshmem_barrier_all_on_stream(streams[stream_id]);
|
||||
|
||||
// Record stop event for this stream on last iteration
|
||||
if (i == loop - 1) {
|
||||
CHECK_HIP(hipEventRecord(stop_events_timed[stream_id],
|
||||
streams[stream_id]));
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
num_msgs = (loop + args.skip) * num_streams;
|
||||
num_timed_msgs = loop * num_streams;
|
||||
}
|
||||
|
||||
void BarrierAllOnStreamTester::verifyResults(size_t size) {}
|
||||
@@ -0,0 +1,67 @@
|
||||
/******************************************************************************
|
||||
* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* SPDX-License-Identifier: MIT
|
||||
*
|
||||
* 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 _BARRIER_ALL_ON_STREAM_TESTER_HPP_
|
||||
#define _BARRIER_ALL_ON_STREAM_TESTER_HPP_
|
||||
|
||||
#include "tester.hpp"
|
||||
#include <vector>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
using namespace rocshmem;
|
||||
|
||||
/******************************************************************************
|
||||
* HOST TESTER CLASS
|
||||
*****************************************************************************/
|
||||
class BarrierAllOnStreamTester : public Tester {
|
||||
public:
|
||||
explicit BarrierAllOnStreamTester(TesterArguments args);
|
||||
virtual ~BarrierAllOnStreamTester();
|
||||
|
||||
protected:
|
||||
virtual void resetBuffers(size_t size) override;
|
||||
|
||||
virtual void preLaunchKernel() override;
|
||||
|
||||
virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
|
||||
size_t size) override;
|
||||
|
||||
virtual void postLaunchKernel() override;
|
||||
|
||||
virtual void verifyResults(size_t size) override;
|
||||
|
||||
private:
|
||||
int my_pe;
|
||||
int n_pes;
|
||||
int num_streams = 1;
|
||||
bool use_default_stream = false;
|
||||
std::vector<hipStream_t> streams;
|
||||
std::vector<hipEvent_t> start_events_timed;
|
||||
std::vector<hipEvent_t> stop_events_timed;
|
||||
};
|
||||
|
||||
#include "barrier_all_on_stream_tester.cpp"
|
||||
|
||||
#endif
|
||||
|
||||
@@ -0,0 +1,205 @@
|
||||
/******************************************************************************
|
||||
* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* SPDX-License-Identifier: MIT
|
||||
*
|
||||
* 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 "getmem_on_stream_tester.hpp"
|
||||
|
||||
#include <rocshmem/rocshmem.hpp>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <cstring>
|
||||
#include <cassert>
|
||||
#include <vector>
|
||||
|
||||
/******************************************************************************
|
||||
* HOST TESTER CLASS METHODS
|
||||
*****************************************************************************/
|
||||
GetmemOnStreamTester::GetmemOnStreamTester(TesterArguments args)
|
||||
: Tester(args) {
|
||||
my_pe = rocshmem_my_pe();
|
||||
n_pes = rocshmem_n_pes();
|
||||
|
||||
char *value{nullptr};
|
||||
if ((value = getenv("ROCSHMEM_TEST_NUM_STREAMS"))) {
|
||||
num_streams = atoi(value);
|
||||
} else {
|
||||
// Default to 1 stream
|
||||
num_streams = 1;
|
||||
}
|
||||
|
||||
// Set target PE to get from (default: next PE in ring)
|
||||
pe_target = (my_pe + 1) % n_pes;
|
||||
if ((value = getenv("ROCSHMEM_TEST_GETMEM_TARGET"))) {
|
||||
pe_target = atoi(value);
|
||||
if (pe_target < 0 || pe_target >= n_pes) {
|
||||
std::cerr << "Invalid ROCSHMEM_TEST_GETMEM_TARGET value. Using next PE."
|
||||
<< std::endl;
|
||||
pe_target = (my_pe + 1) % n_pes;
|
||||
}
|
||||
}
|
||||
|
||||
int num_bytes_stream = args.max_msg_size;
|
||||
int total_bytes = num_bytes_stream * num_streams;
|
||||
buf_size = total_bytes;
|
||||
|
||||
source_buf = static_cast<char *>(rocshmem_malloc(buf_size));
|
||||
dest_buf = static_cast<char *>(rocshmem_malloc(buf_size));
|
||||
|
||||
if (source_buf == nullptr || dest_buf == nullptr) {
|
||||
std::cerr << "Error allocating memory from symmetric heap" << std::endl;
|
||||
std::cerr << "source: " << source_buf << ", dest: " << dest_buf
|
||||
<< std::endl;
|
||||
rocshmem_global_exit(1);
|
||||
}
|
||||
|
||||
streams.resize(num_streams);
|
||||
start_events_timed.resize(num_streams);
|
||||
stop_events_timed.resize(num_streams);
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipStreamCreate(&streams[i]));
|
||||
CHECK_HIP(hipEventCreate(&start_events_timed[i]));
|
||||
CHECK_HIP(hipEventCreate(&stop_events_timed[i]));
|
||||
}
|
||||
}
|
||||
|
||||
GetmemOnStreamTester::~GetmemOnStreamTester() {
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipEventDestroy(stop_events_timed[i]));
|
||||
CHECK_HIP(hipEventDestroy(start_events_timed[i]));
|
||||
CHECK_HIP(hipStreamDestroy(streams[i]));
|
||||
}
|
||||
rocshmem_free(source_buf);
|
||||
rocshmem_free(dest_buf);
|
||||
}
|
||||
|
||||
void GetmemOnStreamTester::preLaunchKernel() {
|
||||
bw_factor = 1; // Point-to-point operation
|
||||
}
|
||||
|
||||
void GetmemOnStreamTester::postLaunchKernel() {
|
||||
// Synchronize all streams to ensure events are recorded
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipStreamSynchronize(streams[i]));
|
||||
}
|
||||
|
||||
// Get elapsed time for each stream from HIP events
|
||||
for (int stream_id = 0; stream_id < num_streams && stream_id < num_timers;
|
||||
stream_id++) {
|
||||
float elapsed_time_ms = 0.0f;
|
||||
CHECK_HIP(hipEventElapsedTime(&elapsed_time_ms,
|
||||
start_events_timed[stream_id],
|
||||
stop_events_timed[stream_id]));
|
||||
|
||||
// Convert milliseconds to GPU cycles
|
||||
// wall_clk_rate is in kHz, so: cycles = ms * wall_clk_rate
|
||||
long long int elapsed_cycles =
|
||||
static_cast<long long int>(elapsed_time_ms *
|
||||
static_cast<float>(wall_clk_rate));
|
||||
|
||||
start_time[stream_id] = 0;
|
||||
end_time[stream_id] = elapsed_cycles;
|
||||
}
|
||||
|
||||
// Fill remaining timers with zero if num_timers > num_streams
|
||||
for (int i = num_streams; i < num_timers; i++) {
|
||||
start_time[i] = 0;
|
||||
end_time[i] = 0;
|
||||
}
|
||||
}
|
||||
|
||||
void GetmemOnStreamTester::resetBuffers(size_t size) {
|
||||
// Initialize source buffer on all PEs
|
||||
// Each stream has its own portion
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
int idx = stream_id * size;
|
||||
// Each PE fills its source buffer with a unique value
|
||||
int value = (my_pe + 1) * 100 + stream_id;
|
||||
std::memset(source_buf + idx, value, size);
|
||||
}
|
||||
|
||||
// Clear destination buffer
|
||||
std::memset(dest_buf, 0, buf_size);
|
||||
}
|
||||
|
||||
void GetmemOnStreamTester::launchKernel(dim3 gridSize, dim3 blockSize,
|
||||
int loop, size_t size) {
|
||||
// Execute warmup iterations (skip)
|
||||
for (int i = 0; i < args.skip; i++) {
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
char *stream_dest = dest_buf + stream_id * size;
|
||||
char *stream_source = source_buf + stream_id * size;
|
||||
rocshmem_getmem_on_stream(stream_dest, stream_source, size, pe_target,
|
||||
streams[stream_id]);
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipStreamSynchronize(streams[i]));
|
||||
}
|
||||
|
||||
for (int i = 0; i < loop; i++) {
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
// Record start event for this stream on first iteration
|
||||
if (i == 0) {
|
||||
CHECK_HIP(hipEventRecord(start_events_timed[stream_id],
|
||||
streams[stream_id]));
|
||||
}
|
||||
|
||||
char *stream_dest = dest_buf + stream_id * size;
|
||||
char *stream_source = source_buf + stream_id * size;
|
||||
rocshmem_getmem_on_stream(stream_dest, stream_source, size, pe_target,
|
||||
streams[stream_id]);
|
||||
|
||||
// Record stop event for this stream on last iteration
|
||||
if (i == loop - 1) {
|
||||
CHECK_HIP(hipEventRecord(stop_events_timed[stream_id],
|
||||
streams[stream_id]));
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
num_msgs = (loop + args.skip) * num_streams;
|
||||
num_timed_msgs = loop * num_streams;
|
||||
}
|
||||
|
||||
void GetmemOnStreamTester::verifyResults(size_t size) {
|
||||
// Verify correctness: after getmem, local dest buffer should have
|
||||
// the data from target PE's source buffer
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
int idx = stream_id * size;
|
||||
// Expected value is from pe_target
|
||||
int expected_value = (pe_target + 1) * 100 + stream_id;
|
||||
|
||||
for (size_t k = 0; k < size; k++) {
|
||||
if (static_cast<unsigned char>(dest_buf[idx + k]) !=
|
||||
static_cast<unsigned char>(expected_value)) {
|
||||
std::cerr << "PE " << my_pe << ": Verification failed for stream "
|
||||
<< stream_id << " at byte " << k << std::endl;
|
||||
std::cerr << "Expected value: " << expected_value
|
||||
<< ", Got: " << static_cast<int>(dest_buf[idx + k])
|
||||
<< std::endl;
|
||||
rocshmem_global_exit(1);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,70 @@
|
||||
/******************************************************************************
|
||||
* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* SPDX-License-Identifier: MIT
|
||||
*
|
||||
* 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 _GETMEM_ON_STREAM_TESTER_HPP_
|
||||
#define _GETMEM_ON_STREAM_TESTER_HPP_
|
||||
|
||||
#include "tester.hpp"
|
||||
#include <vector>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
using namespace rocshmem;
|
||||
|
||||
/******************************************************************************
|
||||
* HOST TESTER CLASS
|
||||
*****************************************************************************/
|
||||
class GetmemOnStreamTester : public Tester {
|
||||
public:
|
||||
explicit GetmemOnStreamTester(TesterArguments args);
|
||||
virtual ~GetmemOnStreamTester();
|
||||
|
||||
protected:
|
||||
virtual void resetBuffers(size_t size) override;
|
||||
|
||||
virtual void preLaunchKernel() override;
|
||||
|
||||
virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
|
||||
size_t size) override;
|
||||
|
||||
virtual void postLaunchKernel() override;
|
||||
|
||||
virtual void verifyResults(size_t size) override;
|
||||
|
||||
private:
|
||||
char *source_buf;
|
||||
char *dest_buf;
|
||||
int my_pe;
|
||||
int n_pes;
|
||||
size_t buf_size;
|
||||
int num_streams = 1;
|
||||
int pe_target; // Target PE to get from
|
||||
std::vector<hipStream_t> streams;
|
||||
std::vector<hipEvent_t> start_events_timed;
|
||||
std::vector<hipEvent_t> stop_events_timed;
|
||||
};
|
||||
|
||||
#include "getmem_on_stream_tester.cpp"
|
||||
|
||||
#endif
|
||||
|
||||
@@ -0,0 +1,225 @@
|
||||
/******************************************************************************
|
||||
* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* SPDX-License-Identifier: MIT
|
||||
*
|
||||
* 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 "putmem_on_stream_tester.hpp"
|
||||
|
||||
#include <rocshmem/rocshmem.hpp>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <cstring>
|
||||
#include <cassert>
|
||||
#include <vector>
|
||||
|
||||
/******************************************************************************
|
||||
* HOST TESTER CLASS METHODS
|
||||
*****************************************************************************/
|
||||
PutmemOnStreamTester::PutmemOnStreamTester(TesterArguments args)
|
||||
: Tester(args) {
|
||||
my_pe = rocshmem_my_pe();
|
||||
n_pes = rocshmem_n_pes();
|
||||
|
||||
char *value{nullptr};
|
||||
if ((value = getenv("ROCSHMEM_TEST_NUM_STREAMS"))) {
|
||||
num_streams = atoi(value);
|
||||
} else {
|
||||
// Default to 1 stream
|
||||
num_streams = 1;
|
||||
}
|
||||
|
||||
// Check if we should test with nullptr (default stream)
|
||||
use_default_stream = false;
|
||||
if ((value = getenv("ROCSHMEM_TEST_USE_DEFAULT_STREAM"))) {
|
||||
use_default_stream = (atoi(value) != 0);
|
||||
if (use_default_stream) {
|
||||
num_streams = 1; // Only test with one nullptr stream
|
||||
}
|
||||
}
|
||||
|
||||
// Set target PE to put to (default: next PE in ring)
|
||||
pe_target = (my_pe + 1) % n_pes;
|
||||
if ((value = getenv("ROCSHMEM_TEST_PUTMEM_TARGET"))) {
|
||||
pe_target = atoi(value);
|
||||
if (pe_target < 0 || pe_target >= n_pes) {
|
||||
std::cerr << "Invalid ROCSHMEM_TEST_PUTMEM_TARGET value. Using next PE."
|
||||
<< std::endl;
|
||||
pe_target = (my_pe + 1) % n_pes;
|
||||
}
|
||||
}
|
||||
|
||||
int num_bytes_stream = args.max_msg_size;
|
||||
int total_bytes = num_bytes_stream * num_streams;
|
||||
buf_size = total_bytes;
|
||||
|
||||
source_buf = static_cast<char *>(rocshmem_malloc(buf_size));
|
||||
dest_buf = static_cast<char *>(rocshmem_malloc(buf_size));
|
||||
|
||||
if (source_buf == nullptr || dest_buf == nullptr) {
|
||||
std::cerr << "Error allocating memory from symmetric heap" << std::endl;
|
||||
std::cerr << "source: " << source_buf << ", dest: " << dest_buf
|
||||
<< std::endl;
|
||||
rocshmem_global_exit(1);
|
||||
}
|
||||
|
||||
streams.resize(num_streams);
|
||||
start_events_timed.resize(num_streams);
|
||||
stop_events_timed.resize(num_streams);
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
if (use_default_stream) {
|
||||
streams[i] = nullptr; // Use default stream (0)
|
||||
} else {
|
||||
CHECK_HIP(hipStreamCreate(&streams[i]));
|
||||
}
|
||||
CHECK_HIP(hipEventCreate(&start_events_timed[i]));
|
||||
CHECK_HIP(hipEventCreate(&stop_events_timed[i]));
|
||||
}
|
||||
}
|
||||
|
||||
PutmemOnStreamTester::~PutmemOnStreamTester() {
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipEventDestroy(stop_events_timed[i]));
|
||||
CHECK_HIP(hipEventDestroy(start_events_timed[i]));
|
||||
// Don't destroy default stream (nullptr)
|
||||
if (!use_default_stream) {
|
||||
CHECK_HIP(hipStreamDestroy(streams[i]));
|
||||
}
|
||||
}
|
||||
rocshmem_free(source_buf);
|
||||
rocshmem_free(dest_buf);
|
||||
}
|
||||
|
||||
void PutmemOnStreamTester::preLaunchKernel() {
|
||||
bw_factor = 1; // Point-to-point operation
|
||||
}
|
||||
|
||||
void PutmemOnStreamTester::postLaunchKernel() {
|
||||
// Synchronize all streams to ensure events are recorded
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipStreamSynchronize(streams[i]));
|
||||
}
|
||||
|
||||
// Get elapsed time for each stream from HIP events
|
||||
for (int stream_id = 0; stream_id < num_streams && stream_id < num_timers;
|
||||
stream_id++) {
|
||||
float elapsed_time_ms = 0.0f;
|
||||
CHECK_HIP(hipEventElapsedTime(&elapsed_time_ms,
|
||||
start_events_timed[stream_id],
|
||||
stop_events_timed[stream_id]));
|
||||
|
||||
// Convert milliseconds to GPU cycles
|
||||
// wall_clk_rate is in kHz, so: cycles = ms * wall_clk_rate
|
||||
long long int elapsed_cycles =
|
||||
static_cast<long long int>(elapsed_time_ms *
|
||||
static_cast<float>(wall_clk_rate));
|
||||
|
||||
start_time[stream_id] = 0;
|
||||
end_time[stream_id] = elapsed_cycles;
|
||||
}
|
||||
|
||||
// Fill remaining timers with zero if num_timers > num_streams
|
||||
for (int i = num_streams; i < num_timers; i++) {
|
||||
start_time[i] = 0;
|
||||
end_time[i] = 0;
|
||||
}
|
||||
}
|
||||
|
||||
void PutmemOnStreamTester::resetBuffers(size_t size) {
|
||||
// Initialize source buffer on all PEs
|
||||
// Each stream has its own portion
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
int idx = stream_id * size;
|
||||
// Each PE fills its source buffer with a unique value
|
||||
int value = (my_pe + 1) * 100 + stream_id;
|
||||
std::memset(source_buf + idx, value, size);
|
||||
}
|
||||
|
||||
// Clear destination buffer (will receive data from other PEs)
|
||||
std::memset(dest_buf, 0, buf_size);
|
||||
}
|
||||
|
||||
void PutmemOnStreamTester::launchKernel(dim3 gridSize, dim3 blockSize,
|
||||
int loop, size_t size) {
|
||||
// Execute warmup iterations (skip)
|
||||
for (int i = 0; i < args.skip; i++) {
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
char *stream_source = source_buf + stream_id * size;
|
||||
char *stream_dest = dest_buf + stream_id * size;
|
||||
rocshmem_putmem_on_stream(stream_dest, stream_source, size, pe_target,
|
||||
streams[stream_id]);
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipStreamSynchronize(streams[i]));
|
||||
}
|
||||
|
||||
for (int i = 0; i < loop; i++) {
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
// Record start event for this stream on first iteration
|
||||
if (i == 0) {
|
||||
CHECK_HIP(hipEventRecord(start_events_timed[stream_id],
|
||||
streams[stream_id]));
|
||||
}
|
||||
|
||||
char *stream_source = source_buf + stream_id * size;
|
||||
char *stream_dest = dest_buf + stream_id * size;
|
||||
rocshmem_putmem_on_stream(stream_dest, stream_source, size, pe_target,
|
||||
streams[stream_id]);
|
||||
|
||||
// Record stop event for this stream on last iteration
|
||||
if (i == loop - 1) {
|
||||
CHECK_HIP(hipEventRecord(stop_events_timed[stream_id],
|
||||
streams[stream_id]));
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
num_msgs = (loop + args.skip) * num_streams;
|
||||
num_timed_msgs = loop * num_streams;
|
||||
}
|
||||
|
||||
void PutmemOnStreamTester::verifyResults(size_t size) {
|
||||
// Verify correctness: after putmem, my dest buffer should have
|
||||
// the data that was put from the PE that targets me
|
||||
// We need to find which PE writes to me: pe_source where (pe_source + 1) % n_pes == my_pe
|
||||
int pe_source = (my_pe - 1 + n_pes) % n_pes;
|
||||
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
int idx = stream_id * size;
|
||||
// Expected value is from pe_source
|
||||
int expected_value = (pe_source + 1) * 100 + stream_id;
|
||||
|
||||
for (size_t k = 0; k < size; k++) {
|
||||
if (static_cast<unsigned char>(dest_buf[idx + k]) !=
|
||||
static_cast<unsigned char>(expected_value)) {
|
||||
std::cerr << "PE " << my_pe << ": Verification failed for stream "
|
||||
<< stream_id << " at byte " << k << std::endl;
|
||||
std::cerr << "Expected value from PE " << pe_source << ": "
|
||||
<< expected_value
|
||||
<< ", Got: " << static_cast<int>(dest_buf[idx + k])
|
||||
<< std::endl;
|
||||
rocshmem_global_exit(1);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,71 @@
|
||||
/******************************************************************************
|
||||
* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* SPDX-License-Identifier: MIT
|
||||
*
|
||||
* 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 _PUTMEM_ON_STREAM_TESTER_HPP_
|
||||
#define _PUTMEM_ON_STREAM_TESTER_HPP_
|
||||
|
||||
#include "tester.hpp"
|
||||
#include <vector>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
using namespace rocshmem;
|
||||
|
||||
/******************************************************************************
|
||||
* HOST TESTER CLASS
|
||||
*****************************************************************************/
|
||||
class PutmemOnStreamTester : public Tester {
|
||||
public:
|
||||
explicit PutmemOnStreamTester(TesterArguments args);
|
||||
virtual ~PutmemOnStreamTester();
|
||||
|
||||
protected:
|
||||
virtual void resetBuffers(size_t size) override;
|
||||
|
||||
virtual void preLaunchKernel() override;
|
||||
|
||||
virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
|
||||
size_t size) override;
|
||||
|
||||
virtual void postLaunchKernel() override;
|
||||
|
||||
virtual void verifyResults(size_t size) override;
|
||||
|
||||
private:
|
||||
char *source_buf;
|
||||
char *dest_buf;
|
||||
int my_pe;
|
||||
int n_pes;
|
||||
size_t buf_size;
|
||||
int num_streams = 1;
|
||||
bool use_default_stream = false;
|
||||
int pe_target; // Target PE to put to
|
||||
std::vector<hipStream_t> streams;
|
||||
std::vector<hipEvent_t> start_events_timed;
|
||||
std::vector<hipEvent_t> stop_events_timed;
|
||||
};
|
||||
|
||||
#include "putmem_on_stream_tester.cpp"
|
||||
|
||||
#endif
|
||||
|
||||
@@ -0,0 +1,236 @@
|
||||
/******************************************************************************
|
||||
* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* SPDX-License-Identifier: MIT
|
||||
*
|
||||
* 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 "putmem_signal_on_stream_tester.hpp"
|
||||
|
||||
#include <rocshmem/rocshmem.hpp>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <cstring>
|
||||
#include <cassert>
|
||||
#include <vector>
|
||||
|
||||
/******************************************************************************
|
||||
* HOST TESTER CLASS METHODS
|
||||
*****************************************************************************/
|
||||
PutmemSignalOnStreamTester::PutmemSignalOnStreamTester(TesterArguments args)
|
||||
: Tester(args) {
|
||||
my_pe = rocshmem_my_pe();
|
||||
n_pes = rocshmem_n_pes();
|
||||
|
||||
char *value{nullptr};
|
||||
if ((value = getenv("ROCSHMEM_TEST_NUM_STREAMS"))) {
|
||||
num_streams = atoi(value);
|
||||
} else {
|
||||
// Default to 1 stream
|
||||
num_streams = 1;
|
||||
}
|
||||
|
||||
// Set target PE to put to (default: next PE in ring)
|
||||
pe_target = (my_pe + 1) % n_pes;
|
||||
if ((value = getenv("ROCSHMEM_TEST_PUTMEM_TARGET"))) {
|
||||
pe_target = atoi(value);
|
||||
if (pe_target < 0 || pe_target >= n_pes) {
|
||||
std::cerr << "Invalid ROCSHMEM_TEST_PUTMEM_TARGET value. Using next PE."
|
||||
<< std::endl;
|
||||
pe_target = (my_pe + 1) % n_pes;
|
||||
}
|
||||
}
|
||||
|
||||
int num_bytes_stream = args.max_msg_size;
|
||||
int total_bytes = num_bytes_stream * num_streams;
|
||||
buf_size = total_bytes;
|
||||
|
||||
source_buf = static_cast<char *>(rocshmem_malloc(buf_size));
|
||||
dest_buf = static_cast<char *>(rocshmem_malloc(buf_size));
|
||||
sig_addr = static_cast<uint64_t *>(rocshmem_malloc(num_streams * sizeof(uint64_t)));
|
||||
|
||||
if (source_buf == nullptr || dest_buf == nullptr || sig_addr == nullptr) {
|
||||
std::cerr << "Error allocating memory from symmetric heap" << std::endl;
|
||||
std::cerr << "source: " << source_buf << ", dest: " << dest_buf
|
||||
<< ", sig_addr: " << sig_addr << std::endl;
|
||||
rocshmem_global_exit(1);
|
||||
}
|
||||
|
||||
streams.resize(num_streams);
|
||||
start_events_timed.resize(num_streams);
|
||||
stop_events_timed.resize(num_streams);
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipStreamCreate(&streams[i]));
|
||||
CHECK_HIP(hipEventCreate(&start_events_timed[i]));
|
||||
CHECK_HIP(hipEventCreate(&stop_events_timed[i]));
|
||||
}
|
||||
}
|
||||
|
||||
PutmemSignalOnStreamTester::~PutmemSignalOnStreamTester() {
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipEventDestroy(stop_events_timed[i]));
|
||||
CHECK_HIP(hipEventDestroy(start_events_timed[i]));
|
||||
CHECK_HIP(hipStreamDestroy(streams[i]));
|
||||
}
|
||||
rocshmem_free(source_buf);
|
||||
rocshmem_free(dest_buf);
|
||||
rocshmem_free(sig_addr);
|
||||
}
|
||||
|
||||
void PutmemSignalOnStreamTester::preLaunchKernel() {
|
||||
bw_factor = 1; // Point-to-point operation
|
||||
}
|
||||
|
||||
void PutmemSignalOnStreamTester::postLaunchKernel() {
|
||||
// Synchronize all streams to ensure events are recorded
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipStreamSynchronize(streams[i]));
|
||||
}
|
||||
|
||||
// Get elapsed time for each stream from HIP events
|
||||
for (int stream_id = 0; stream_id < num_streams && stream_id < num_timers;
|
||||
stream_id++) {
|
||||
float elapsed_time_ms = 0.0f;
|
||||
CHECK_HIP(hipEventElapsedTime(&elapsed_time_ms,
|
||||
start_events_timed[stream_id],
|
||||
stop_events_timed[stream_id]));
|
||||
|
||||
// Convert milliseconds to GPU cycles
|
||||
// wall_clk_rate is in kHz, so: cycles = ms * wall_clk_rate
|
||||
long long int elapsed_cycles =
|
||||
static_cast<long long int>(elapsed_time_ms *
|
||||
static_cast<float>(wall_clk_rate));
|
||||
|
||||
start_time[stream_id] = 0;
|
||||
end_time[stream_id] = elapsed_cycles;
|
||||
}
|
||||
|
||||
// Fill remaining timers with zero if num_timers > num_streams
|
||||
for (int i = num_streams; i < num_timers; i++) {
|
||||
start_time[i] = 0;
|
||||
end_time[i] = 0;
|
||||
}
|
||||
}
|
||||
|
||||
void PutmemSignalOnStreamTester::resetBuffers(size_t size) {
|
||||
// Initialize source buffer on all PEs
|
||||
// Each stream has its own portion
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
int idx = stream_id * size;
|
||||
// Each PE fills its source buffer with a unique value
|
||||
int value = (my_pe + 1) * 100 + stream_id;
|
||||
std::memset(source_buf + idx, value, size);
|
||||
}
|
||||
|
||||
// Clear destination buffer (will receive data from other PEs)
|
||||
std::memset(dest_buf, 0, buf_size);
|
||||
|
||||
// Clear signal addresses
|
||||
std::memset(sig_addr, 0, num_streams * sizeof(uint64_t));
|
||||
}
|
||||
|
||||
void PutmemSignalOnStreamTester::launchKernel(dim3 gridSize, dim3 blockSize,
|
||||
int loop, size_t size) {
|
||||
uint64_t signal_value = 1;
|
||||
|
||||
// Execute warmup iterations (skip)
|
||||
for (int i = 0; i < args.skip; i++) {
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
char *stream_source = source_buf + stream_id * size;
|
||||
char *stream_dest = dest_buf + stream_id * size;
|
||||
rocshmem_putmem_signal_on_stream(stream_dest, stream_source, size,
|
||||
&sig_addr[stream_id], signal_value,
|
||||
sig_op, pe_target, streams[stream_id]);
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipStreamSynchronize(streams[i]));
|
||||
}
|
||||
|
||||
// Reset signal addresses after warmup and synchronize across PEs
|
||||
std::memset(sig_addr, 0, num_streams * sizeof(uint64_t));
|
||||
rocshmem_barrier_all();
|
||||
|
||||
for (int i = 0; i < loop; i++) {
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
// Record start event for this stream on first iteration
|
||||
if (i == 0) {
|
||||
CHECK_HIP(hipEventRecord(start_events_timed[stream_id],
|
||||
streams[stream_id]));
|
||||
}
|
||||
|
||||
char *stream_source = source_buf + stream_id * size;
|
||||
char *stream_dest = dest_buf + stream_id * size;
|
||||
rocshmem_putmem_signal_on_stream(stream_dest, stream_source, size,
|
||||
&sig_addr[stream_id], signal_value,
|
||||
sig_op, pe_target, streams[stream_id]);
|
||||
|
||||
// Record stop event for this stream on last iteration
|
||||
if (i == loop - 1) {
|
||||
CHECK_HIP(hipEventRecord(stop_events_timed[stream_id],
|
||||
streams[stream_id]));
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
num_msgs = (loop + args.skip) * num_streams;
|
||||
num_timed_msgs = loop * num_streams;
|
||||
}
|
||||
|
||||
void PutmemSignalOnStreamTester::verifyResults(size_t size) {
|
||||
// Synchronize to ensure all operations completed
|
||||
rocshmem_barrier_all();
|
||||
|
||||
// Verify correctness: after putmem_signal, my dest buffer should have
|
||||
// the data that was put from the PE that targets me
|
||||
// We need to find which PE writes to me: pe_source where (pe_source + 1) % n_pes == my_pe
|
||||
int pe_source = (my_pe - 1 + n_pes) % n_pes;
|
||||
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
int idx = stream_id * size;
|
||||
// Expected value is from pe_source
|
||||
int expected_value = (pe_source + 1) * 100 + stream_id;
|
||||
|
||||
// Verify data
|
||||
for (size_t k = 0; k < size; k++) {
|
||||
if (static_cast<unsigned char>(dest_buf[idx + k]) !=
|
||||
static_cast<unsigned char>(expected_value)) {
|
||||
std::cerr << "PE " << my_pe << ": Data verification failed for stream "
|
||||
<< stream_id << " at byte " << k << std::endl;
|
||||
std::cerr << "Expected value from PE " << pe_source << ": "
|
||||
<< expected_value
|
||||
<< ", Got: " << static_cast<int>(dest_buf[idx + k])
|
||||
<< std::endl;
|
||||
rocshmem_global_exit(1);
|
||||
}
|
||||
}
|
||||
|
||||
// Verify signal
|
||||
uint64_t expected_signal = 1;
|
||||
if (sig_addr[stream_id] != expected_signal) {
|
||||
std::cerr << "PE " << my_pe << ": Signal verification failed for stream "
|
||||
<< stream_id << std::endl;
|
||||
std::cerr << "Expected signal: " << expected_signal
|
||||
<< ", Got: " << sig_addr[stream_id] << std::endl;
|
||||
rocshmem_global_exit(1);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,72 @@
|
||||
/******************************************************************************
|
||||
* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* SPDX-License-Identifier: MIT
|
||||
*
|
||||
* 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 _PUTMEM_SIGNAL_ON_STREAM_TESTER_HPP_
|
||||
#define _PUTMEM_SIGNAL_ON_STREAM_TESTER_HPP_
|
||||
|
||||
#include "tester.hpp"
|
||||
#include <vector>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
using namespace rocshmem;
|
||||
|
||||
/******************************************************************************
|
||||
* HOST TESTER CLASS
|
||||
*****************************************************************************/
|
||||
class PutmemSignalOnStreamTester : public Tester {
|
||||
public:
|
||||
explicit PutmemSignalOnStreamTester(TesterArguments args);
|
||||
virtual ~PutmemSignalOnStreamTester();
|
||||
|
||||
protected:
|
||||
virtual void resetBuffers(size_t size) override;
|
||||
|
||||
virtual void preLaunchKernel() override;
|
||||
|
||||
virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
|
||||
size_t size) override;
|
||||
|
||||
virtual void postLaunchKernel() override;
|
||||
|
||||
virtual void verifyResults(size_t size) override;
|
||||
|
||||
private:
|
||||
char *source_buf;
|
||||
char *dest_buf;
|
||||
uint64_t *sig_addr;
|
||||
int my_pe;
|
||||
int n_pes;
|
||||
size_t buf_size;
|
||||
int num_streams = 1;
|
||||
int pe_target; // Target PE to put to
|
||||
int sig_op = ROCSHMEM_SIGNAL_SET; // Signal operation
|
||||
std::vector<hipStream_t> streams;
|
||||
std::vector<hipEvent_t> start_events_timed;
|
||||
std::vector<hipEvent_t> stop_events_timed;
|
||||
};
|
||||
|
||||
#include "putmem_signal_on_stream_tester.cpp"
|
||||
|
||||
#endif
|
||||
|
||||
@@ -0,0 +1,204 @@
|
||||
/******************************************************************************
|
||||
* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* SPDX-License-Identifier: MIT
|
||||
*
|
||||
* 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 <rocshmem/rocshmem.hpp>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
#include <cstring>
|
||||
#include <cassert>
|
||||
|
||||
/******************************************************************************
|
||||
* HOST TESTER CLASS METHODS
|
||||
*****************************************************************************/
|
||||
SignalWaitUntilOnStreamTester::SignalWaitUntilOnStreamTester(
|
||||
TesterArguments args)
|
||||
: Tester(args) {
|
||||
my_pe = rocshmem_my_pe();
|
||||
n_pes = rocshmem_n_pes();
|
||||
|
||||
char *value{nullptr};
|
||||
if ((value = getenv("ROCSHMEM_TEST_NUM_STREAMS"))) {
|
||||
num_streams = atoi(value);
|
||||
} else {
|
||||
// Default to 1 stream
|
||||
num_streams = 1;
|
||||
}
|
||||
|
||||
// Set target PE (next PE in ring)
|
||||
pe_target = (my_pe + 1) % n_pes;
|
||||
|
||||
// Allocate signal addresses on symmetric heap
|
||||
sig_addr =
|
||||
static_cast<uint64_t *>(rocshmem_malloc(num_streams * sizeof(uint64_t)));
|
||||
source_buf =
|
||||
static_cast<uint64_t *>(rocshmem_malloc(num_streams * sizeof(uint64_t)));
|
||||
|
||||
if (sig_addr == nullptr || source_buf == nullptr) {
|
||||
std::cerr << "Error allocating memory from symmetric heap" << std::endl;
|
||||
std::cerr << "sig_addr: " << sig_addr << ", source_buf: " << source_buf
|
||||
<< std::endl;
|
||||
rocshmem_global_exit(1);
|
||||
}
|
||||
|
||||
streams.resize(num_streams);
|
||||
start_events_timed.resize(num_streams);
|
||||
stop_events_timed.resize(num_streams);
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipStreamCreate(&streams[i]));
|
||||
CHECK_HIP(hipEventCreate(&start_events_timed[i]));
|
||||
CHECK_HIP(hipEventCreate(&stop_events_timed[i]));
|
||||
}
|
||||
}
|
||||
|
||||
SignalWaitUntilOnStreamTester::~SignalWaitUntilOnStreamTester() {
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipEventDestroy(stop_events_timed[i]));
|
||||
CHECK_HIP(hipEventDestroy(start_events_timed[i]));
|
||||
CHECK_HIP(hipStreamDestroy(streams[i]));
|
||||
}
|
||||
rocshmem_free(sig_addr);
|
||||
rocshmem_free(source_buf);
|
||||
}
|
||||
|
||||
void SignalWaitUntilOnStreamTester::preLaunchKernel() {
|
||||
bw_factor = 1; // Point-to-point operation
|
||||
}
|
||||
|
||||
void SignalWaitUntilOnStreamTester::postLaunchKernel() {
|
||||
// Synchronize all streams to ensure events are recorded
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
CHECK_HIP(hipStreamSynchronize(streams[i]));
|
||||
}
|
||||
|
||||
// Get elapsed time for each stream from HIP events
|
||||
for (int stream_id = 0; stream_id < num_streams && stream_id < num_timers;
|
||||
stream_id++) {
|
||||
float elapsed_time_ms = 0.0f;
|
||||
CHECK_HIP(hipEventElapsedTime(&elapsed_time_ms,
|
||||
start_events_timed[stream_id],
|
||||
stop_events_timed[stream_id]));
|
||||
|
||||
// Convert milliseconds to GPU cycles
|
||||
long long int elapsed_cycles =
|
||||
static_cast<long long int>(elapsed_time_ms *
|
||||
static_cast<float>(wall_clk_rate));
|
||||
|
||||
start_time[stream_id] = 0;
|
||||
end_time[stream_id] = elapsed_cycles;
|
||||
}
|
||||
|
||||
// Fill remaining timers with zero if num_timers > num_streams
|
||||
for (int i = num_streams; i < num_timers; i++) {
|
||||
start_time[i] = 0;
|
||||
end_time[i] = 0;
|
||||
}
|
||||
}
|
||||
|
||||
void SignalWaitUntilOnStreamTester::resetBuffers(size_t size) {
|
||||
// Clear signal addresses
|
||||
std::memset(sig_addr, 0, num_streams * sizeof(uint64_t));
|
||||
}
|
||||
|
||||
void SignalWaitUntilOnStreamTester::launchKernel(dim3 gridSize, dim3 blockSize,
|
||||
int loop, size_t size) {
|
||||
// Execute warmup + timed iterations
|
||||
for (int i = 0; i < args.skip + loop; i++) {
|
||||
// Increment signal value for each iteration
|
||||
uint64_t signal_value = i + 1;
|
||||
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
// Record start event after warmup on first timed iteration for all streams
|
||||
if (i == args.skip) {
|
||||
CHECK_HIP(hipEventRecord(start_events_timed[stream_id],
|
||||
streams[stream_id]));
|
||||
}
|
||||
|
||||
// PE 0 starts the ring by signaling PE 1
|
||||
if (my_pe == 0) {
|
||||
rocshmem_putmem_signal_on_stream(&sig_addr[stream_id],
|
||||
&source_buf[stream_id],
|
||||
sizeof(uint64_t), &sig_addr[stream_id],
|
||||
signal_value, sig_op, pe_target,
|
||||
streams[stream_id]);
|
||||
} else {
|
||||
// All other PEs wait for signal from previous PE
|
||||
rocshmem_signal_wait_until_on_stream(&sig_addr[stream_id],
|
||||
ROCSHMEM_CMP_GE, signal_value,
|
||||
streams[stream_id]);
|
||||
|
||||
// Forward the signal to next PE (unless we're the last PE)
|
||||
if (my_pe != n_pes - 1) {
|
||||
rocshmem_putmem_signal_on_stream(&sig_addr[stream_id],
|
||||
&source_buf[stream_id],
|
||||
sizeof(uint64_t), &sig_addr[stream_id],
|
||||
signal_value, sig_op, pe_target,
|
||||
streams[stream_id]);
|
||||
}
|
||||
}
|
||||
|
||||
// Record stop event on last timed iteration for all streams
|
||||
if (i == args.skip + loop - 1) {
|
||||
CHECK_HIP(hipEventRecord(stop_events_timed[stream_id],
|
||||
streams[stream_id]));
|
||||
}
|
||||
}
|
||||
|
||||
// Wait for all streams to complete
|
||||
for (int j = 0; j < num_streams; j++) {
|
||||
CHECK_HIP(hipStreamSynchronize(streams[j]));
|
||||
}
|
||||
|
||||
// Barrier to ensure all RMA operations completed across all PEs
|
||||
rocshmem_barrier_all();
|
||||
}
|
||||
|
||||
num_msgs = (loop + args.skip) * num_streams;
|
||||
num_timed_msgs = loop * num_streams;
|
||||
}
|
||||
|
||||
void SignalWaitUntilOnStreamTester::verifyResults(size_t size) {
|
||||
// Synchronize to ensure all operations completed
|
||||
rocshmem_barrier_all();
|
||||
|
||||
// Verify signal values
|
||||
// All PEs except PE 0 should have received the final signal value
|
||||
uint64_t expected_signal = args.skip + args.loop;
|
||||
|
||||
for (int stream_id = 0; stream_id < num_streams; stream_id++) {
|
||||
// PE 0 doesn't receive signals (it initiates), so skip verification
|
||||
if (my_pe == 0) {
|
||||
continue;
|
||||
}
|
||||
|
||||
// Verify signal
|
||||
if (sig_addr[stream_id] != expected_signal) {
|
||||
std::cerr << "PE " << my_pe << ": Signal verification failed for stream "
|
||||
<< stream_id << std::endl;
|
||||
std::cerr << "Expected signal: " << expected_signal
|
||||
<< ", Got: " << sig_addr[stream_id] << std::endl;
|
||||
rocshmem_global_exit(1);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,69 @@
|
||||
/******************************************************************************
|
||||
* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* SPDX-License-Identifier: MIT
|
||||
*
|
||||
* 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 _SIGNAL_WAIT_UNTIL_ON_STREAM_TESTER_HPP_
|
||||
#define _SIGNAL_WAIT_UNTIL_ON_STREAM_TESTER_HPP_
|
||||
|
||||
#include "tester.hpp"
|
||||
#include <vector>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
using namespace rocshmem;
|
||||
|
||||
/******************************************************************************
|
||||
* HOST TESTER CLASS
|
||||
*****************************************************************************/
|
||||
class SignalWaitUntilOnStreamTester : public Tester {
|
||||
public:
|
||||
explicit SignalWaitUntilOnStreamTester(TesterArguments args);
|
||||
virtual ~SignalWaitUntilOnStreamTester();
|
||||
|
||||
protected:
|
||||
virtual void resetBuffers(size_t size) override;
|
||||
|
||||
virtual void preLaunchKernel() override;
|
||||
|
||||
virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
|
||||
size_t size) override;
|
||||
|
||||
virtual void postLaunchKernel() override;
|
||||
|
||||
virtual void verifyResults(size_t size) override;
|
||||
|
||||
private:
|
||||
uint64_t *sig_addr;
|
||||
uint64_t *source_buf; // Source buffer in symmetric heap
|
||||
int my_pe;
|
||||
int n_pes;
|
||||
int num_streams = 1;
|
||||
int pe_target; // Target PE to signal next
|
||||
int sig_op = ROCSHMEM_SIGNAL_SET; // Signal operation
|
||||
std::vector<hipStream_t> streams;
|
||||
std::vector<hipEvent_t> start_events_timed;
|
||||
std::vector<hipEvent_t> stop_events_timed;
|
||||
};
|
||||
|
||||
#include "signal_wait_until_on_stream_tester.cpp"
|
||||
|
||||
#endif
|
||||
@@ -39,7 +39,7 @@ TeamAlltoallmemOnStreamTester::TeamAlltoallmemOnStreamTester(TesterArguments arg
|
||||
n_pes = rocshmem_team_n_pes(ROCSHMEM_TEAM_WORLD);
|
||||
|
||||
char* value{nullptr};
|
||||
if ((value = getenv("ROCSHMEM_MAX_NUM_TEAMS"))) {
|
||||
if ((value = getenv("ROCSHMEM_TEST_MAX_NUM_TEAMS"))) {
|
||||
num_teams = atoi(value);
|
||||
} else {
|
||||
// Default to number of work groups
|
||||
@@ -162,6 +162,10 @@ void TeamAlltoallmemOnStreamTester::launchKernel(dim3 gridSize,
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < num_teams; i++) {
|
||||
CHECK_HIP(hipStreamSynchronize(streams[i]));
|
||||
}
|
||||
|
||||
for (int i = 0; i < loop; i++) {
|
||||
for (int wg_id = 0; wg_id < num_teams; wg_id++) {
|
||||
// Record start event for this work group on first iteration
|
||||
|
||||
@@ -0,0 +1,240 @@
|
||||
/******************************************************************************
|
||||
* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* SPDX-License-Identifier: MIT
|
||||
*
|
||||
* 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 "team_broadcastmem_on_stream_tester.hpp"
|
||||
|
||||
#include <rocshmem/rocshmem.hpp>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <cstring>
|
||||
#include <cassert>
|
||||
#include <vector>
|
||||
|
||||
/******************************************************************************
|
||||
* HOST TESTER CLASS METHODS
|
||||
*****************************************************************************/
|
||||
TeamBroadcastmemOnStreamTester::TeamBroadcastmemOnStreamTester(TesterArguments args)
|
||||
: Tester(args) {
|
||||
my_pe = rocshmem_team_my_pe(ROCSHMEM_TEAM_WORLD);
|
||||
n_pes = rocshmem_team_n_pes(ROCSHMEM_TEAM_WORLD);
|
||||
|
||||
char* value{nullptr};
|
||||
if ((value = getenv("ROCSHMEM_TEST_MAX_NUM_TEAMS"))) {
|
||||
num_teams = atoi(value);
|
||||
} else {
|
||||
// Default to number of work groups
|
||||
num_teams = args.num_wgs;
|
||||
}
|
||||
|
||||
// Set root PE to 0 by default, can be modified via environment variable
|
||||
if ((value = getenv("ROCSHMEM_TEST_BROADCAST_ROOT"))) {
|
||||
pe_root = atoi(value);
|
||||
if (pe_root < 0 || pe_root >= n_pes) {
|
||||
std::cerr << "Invalid ROCSHMEM_TEST_BROADCAST_ROOT value. Using PE 0."
|
||||
<< std::endl;
|
||||
pe_root = 0;
|
||||
}
|
||||
}
|
||||
|
||||
int num_bytes_wg = args.max_msg_size;
|
||||
int total_bytes = num_bytes_wg * num_teams;
|
||||
buf_size = total_bytes;
|
||||
|
||||
source_buf = static_cast<char *>(rocshmem_malloc(buf_size));
|
||||
dest_buf = static_cast<char *>(rocshmem_malloc(buf_size));
|
||||
|
||||
if (source_buf == nullptr || dest_buf == nullptr) {
|
||||
std::cerr << "Error allocating memory from symmetric heap" << std::endl;
|
||||
std::cerr << "source: " << source_buf << ", dest: " << dest_buf
|
||||
<< std::endl;
|
||||
rocshmem_global_exit(1);
|
||||
}
|
||||
|
||||
team_world_dup.resize(num_teams);
|
||||
|
||||
streams.resize(num_teams);
|
||||
start_events_timed.resize(num_teams);
|
||||
stop_events_timed.resize(num_teams);
|
||||
for (int i = 0; i < num_teams; i++) {
|
||||
CHECK_HIP(hipStreamCreate(&streams[i]));
|
||||
CHECK_HIP(hipEventCreate(&start_events_timed[i]));
|
||||
CHECK_HIP(hipEventCreate(&stop_events_timed[i]));
|
||||
}
|
||||
}
|
||||
|
||||
TeamBroadcastmemOnStreamTester::~TeamBroadcastmemOnStreamTester() {
|
||||
for (int i = 0; i < num_teams; i++) {
|
||||
CHECK_HIP(hipEventDestroy(stop_events_timed[i]));
|
||||
CHECK_HIP(hipEventDestroy(start_events_timed[i]));
|
||||
CHECK_HIP(hipStreamDestroy(streams[i]));
|
||||
}
|
||||
rocshmem_free(source_buf);
|
||||
rocshmem_free(dest_buf);
|
||||
}
|
||||
|
||||
void TeamBroadcastmemOnStreamTester::preLaunchKernel() {
|
||||
bw_factor = 1; // Broadcast is one-to-all
|
||||
|
||||
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) {
|
||||
std::cerr << "Team " << team_i << " is invalid!" << std::endl;
|
||||
abort();
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
void TeamBroadcastmemOnStreamTester::postLaunchKernel() {
|
||||
// Synchronize all streams to ensure events are recorded
|
||||
for (int i = 0; i < num_teams; i++) {
|
||||
CHECK_HIP(hipStreamSynchronize(streams[i]));
|
||||
}
|
||||
|
||||
// Get elapsed time for each work group from HIP events
|
||||
for (int wg_id = 0; wg_id < num_teams && wg_id < num_timers; wg_id++) {
|
||||
float elapsed_time_ms = 0.0f;
|
||||
CHECK_HIP(hipEventElapsedTime(&elapsed_time_ms, start_events_timed[wg_id],
|
||||
stop_events_timed[wg_id]));
|
||||
|
||||
// Convert milliseconds to GPU cycles
|
||||
// wall_clk_rate is in kHz, so: cycles = ms * wall_clk_rate
|
||||
long long int elapsed_cycles = static_cast<long long int>(
|
||||
elapsed_time_ms * static_cast<float>(wall_clk_rate));
|
||||
|
||||
start_time[wg_id] = 0;
|
||||
end_time[wg_id] = elapsed_cycles;
|
||||
}
|
||||
|
||||
// Fill remaining timers with zero if num_timers > num_teams
|
||||
for (int i = num_teams; i < num_timers; i++) {
|
||||
start_time[i] = 0;
|
||||
end_time[i] = 0;
|
||||
}
|
||||
|
||||
for (int team_i = 0; team_i < num_teams; team_i++) {
|
||||
rocshmem_team_destroy(team_world_dup[team_i]);
|
||||
}
|
||||
}
|
||||
|
||||
void TeamBroadcastmemOnStreamTester::resetBuffers(size_t size) {
|
||||
// Initialize source buffer on all PEs
|
||||
// Each work group has its own portion
|
||||
for (int wg_id = 0; wg_id < num_teams; wg_id++) {
|
||||
int idx = wg_id * size;
|
||||
if (my_pe == pe_root) {
|
||||
// Root PE fills its source buffer with broadcast value
|
||||
int value = (pe_root + 1) * 100 + wg_id;
|
||||
std::memset(source_buf + idx, value, size);
|
||||
} else {
|
||||
// Non-root PEs source buffer (not used in broadcast)
|
||||
std::memset(source_buf + idx, 0xFF, size);
|
||||
}
|
||||
}
|
||||
|
||||
// Initialize destination buffer on all PEs
|
||||
// Root PE keeps its initial dest value (broadcast doesn't copy to root's
|
||||
// dest) Non-root PEs set to 0 (will receive broadcast data)
|
||||
for (int wg_id = 0; wg_id < num_teams; wg_id++) {
|
||||
int idx = wg_id * size;
|
||||
if (my_pe == pe_root) {
|
||||
// Root PE's dest buffer stays with a different value
|
||||
int root_dest_value = 0xAA;
|
||||
std::memset(dest_buf + idx, root_dest_value, size);
|
||||
} else {
|
||||
std::memset(dest_buf + idx, 0, size);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
void TeamBroadcastmemOnStreamTester::launchKernel(dim3 gridSize,
|
||||
dim3 blockSize,
|
||||
int loop,
|
||||
size_t size) {
|
||||
// Execute warmup iterations (skip)
|
||||
for (int i = 0; i < args.skip; i++) {
|
||||
for (int wg_id = 0; wg_id < num_teams; wg_id++) {
|
||||
char *wg_source = source_buf + wg_id * size;
|
||||
char *wg_dest = dest_buf + wg_id * size;
|
||||
rocshmem_broadcastmem_on_stream(team_world_dup[wg_id], wg_dest,
|
||||
wg_source, size, pe_root, streams[wg_id]);
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < num_teams; i++) {
|
||||
CHECK_HIP(hipStreamSynchronize(streams[i]));
|
||||
}
|
||||
|
||||
for (int i = 0; i < loop; i++) {
|
||||
for (int wg_id = 0; wg_id < num_teams; wg_id++) {
|
||||
// Record start event for this work group on first iteration
|
||||
if (i == 0) {
|
||||
CHECK_HIP(hipEventRecord(start_events_timed[wg_id], streams[wg_id]));
|
||||
}
|
||||
|
||||
char *wg_source = source_buf + wg_id * size;
|
||||
char *wg_dest = dest_buf + wg_id * size;
|
||||
rocshmem_broadcastmem_on_stream(team_world_dup[wg_id], wg_dest,
|
||||
wg_source, size, pe_root, streams[wg_id]);
|
||||
|
||||
// Record stop event for this work group on last iteration
|
||||
if (i == loop - 1) {
|
||||
CHECK_HIP(hipEventRecord(stop_events_timed[wg_id], streams[wg_id]));
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
num_msgs = (loop + args.skip) * num_teams;
|
||||
num_timed_msgs = loop * num_teams;
|
||||
}
|
||||
|
||||
void TeamBroadcastmemOnStreamTester::verifyResults(size_t size) {
|
||||
// Verify correctness: after broadcast, non-root PEs receive the broadcast
|
||||
// data Root PE's dest buffer is NOT modified (per OpenSHMEM/rocSHMEM spec)
|
||||
for (int wg_id = 0; wg_id < num_teams; wg_id++) {
|
||||
int idx = wg_id * size;
|
||||
int expected_value;
|
||||
|
||||
if (my_pe == pe_root) {
|
||||
// Root PE's dest buffer should remain unchanged (0xAA)
|
||||
expected_value = 0xAA;
|
||||
} else {
|
||||
// Non-root PEs should have received the broadcast value
|
||||
expected_value = (pe_root + 1) * 100 + wg_id;
|
||||
}
|
||||
|
||||
for (size_t k = 0; k < size; k++) {
|
||||
if (static_cast<unsigned char>(dest_buf[idx + k]) !=
|
||||
static_cast<unsigned char>(expected_value)) {
|
||||
std::cerr << "PE " << my_pe << ": Verification failed for WG "
|
||||
<< wg_id << " at byte " << k << std::endl;
|
||||
std::cerr << "Expected value: " << expected_value
|
||||
<< ", Got: " << static_cast<int>(dest_buf[idx + k])
|
||||
<< std::endl;
|
||||
rocshmem_global_exit(1);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,71 @@
|
||||
/******************************************************************************
|
||||
* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
|
||||
*
|
||||
* SPDX-License-Identifier: MIT
|
||||
*
|
||||
* 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 _TEAM_BROADCASTMEM_ON_STREAM_TESTER_HPP_
|
||||
#define _TEAM_BROADCASTMEM_ON_STREAM_TESTER_HPP_
|
||||
|
||||
#include "tester.hpp"
|
||||
#include <vector>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
using namespace rocshmem;
|
||||
|
||||
/******************************************************************************
|
||||
* HOST TESTER CLASS
|
||||
*****************************************************************************/
|
||||
class TeamBroadcastmemOnStreamTester : public Tester {
|
||||
public:
|
||||
explicit TeamBroadcastmemOnStreamTester(TesterArguments args);
|
||||
virtual ~TeamBroadcastmemOnStreamTester();
|
||||
|
||||
protected:
|
||||
virtual void resetBuffers(size_t size) override;
|
||||
|
||||
virtual void preLaunchKernel() override;
|
||||
|
||||
virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
|
||||
size_t size) override;
|
||||
|
||||
virtual void postLaunchKernel() override;
|
||||
|
||||
virtual void verifyResults(size_t size) override;
|
||||
|
||||
private:
|
||||
char *source_buf;
|
||||
char *dest_buf;
|
||||
int my_pe;
|
||||
int n_pes;
|
||||
size_t buf_size;
|
||||
int num_teams = 1;
|
||||
int pe_root = 0; // Root PE for broadcast
|
||||
std::vector<rocshmem_team_t> team_world_dup;
|
||||
std::vector<hipStream_t> streams;
|
||||
std::vector<hipEvent_t> start_events_timed;
|
||||
std::vector<hipEvent_t> stop_events_timed;
|
||||
};
|
||||
|
||||
#include "team_broadcastmem_on_stream_tester.cpp"
|
||||
|
||||
#endif
|
||||
|
||||
@@ -36,7 +36,12 @@
|
||||
#include "amo_standard_tester.hpp"
|
||||
#include "default_ctx_primitive_tester.hpp"
|
||||
#include "barrier_all_tester.hpp"
|
||||
#include "barrier_all_on_stream_tester.hpp"
|
||||
#include "empty_tester.hpp"
|
||||
#include "getmem_on_stream_tester.hpp"
|
||||
#include "putmem_on_stream_tester.hpp"
|
||||
#include "putmem_signal_on_stream_tester.hpp"
|
||||
#include "signal_wait_until_on_stream_tester.hpp"
|
||||
#include "ping_all_tester.hpp"
|
||||
#include "ping_pong_tester.hpp"
|
||||
#include "primitive_mr_tester.hpp"
|
||||
@@ -48,6 +53,7 @@
|
||||
#include "team_sync_tester.hpp"
|
||||
#include "team_alltoall_tester.hpp"
|
||||
#include "team_alltoallmem_on_stream_tester.hpp"
|
||||
#include "team_broadcastmem_on_stream_tester.hpp"
|
||||
#include "team_barrier_tester.hpp"
|
||||
#include "team_broadcast_tester.hpp"
|
||||
#include "team_ctx_infra_tester.hpp"
|
||||
@@ -233,6 +239,36 @@ std::vector<Tester*> Tester::create(TesterArguments args) {
|
||||
std::cout << "Alltoallmem_On_Stream ###" << std::endl;
|
||||
testers.push_back(new TeamAlltoallmemOnStreamTester(args));
|
||||
return testers;
|
||||
case BarrierAllOnStreamTestType:
|
||||
if (rank == 0)
|
||||
std::cout << "Barrier_All_On_Stream ###" << std::endl;
|
||||
testers.push_back(new BarrierAllOnStreamTester(args));
|
||||
return testers;
|
||||
case TeamBroadcastmemOnStreamTestType:
|
||||
if (rank == 0)
|
||||
std::cout << "Broadcastmem_On_Stream ###" << std::endl;
|
||||
testers.push_back(new TeamBroadcastmemOnStreamTester(args));
|
||||
return testers;
|
||||
case GetmemOnStreamTestType:
|
||||
if (rank == 0)
|
||||
std::cout << "Getmem_On_Stream ###" << std::endl;
|
||||
testers.push_back(new GetmemOnStreamTester(args));
|
||||
return testers;
|
||||
case PutmemOnStreamTestType:
|
||||
if (rank == 0)
|
||||
std::cout << "Putmem_On_Stream ###" << std::endl;
|
||||
testers.push_back(new PutmemOnStreamTester(args));
|
||||
return testers;
|
||||
case PutmemSignalOnStreamTestType:
|
||||
if (rank == 0)
|
||||
std::cout << "Putmem_Signal_On_Stream ###" << std::endl;
|
||||
testers.push_back(new PutmemSignalOnStreamTester(args));
|
||||
return testers;
|
||||
case SignalWaitUntilOnStreamTestType:
|
||||
if (rank == 0)
|
||||
std::cout << "Signal_Wait_Until_On_Stream ###" << std::endl;
|
||||
testers.push_back(new SignalWaitUntilOnStreamTester(args));
|
||||
return testers;
|
||||
case TeamFCollectTestType:
|
||||
if (rank == 0) {
|
||||
std::cout << "Fcollect Test ###" << std::endl;
|
||||
@@ -569,30 +605,50 @@ void Tester::execute() {
|
||||
}
|
||||
|
||||
bool Tester::peLaunchesKernel() {
|
||||
bool is_launcher;
|
||||
|
||||
/**
|
||||
* The PE assigned 0 is always active in these tests.
|
||||
*/
|
||||
is_launcher = args.myid == 0;
|
||||
bool is_launcher = (args.myid == 0);
|
||||
|
||||
/**
|
||||
* Some test types are active on both sides.
|
||||
*/
|
||||
is_launcher = is_launcher || (_type == TeamReductionTestType) ||
|
||||
(_type == TeamBroadcastTestType) || (_type == TeamCtxInfraTestType) ||
|
||||
(_type == TeamCtxInfraTestSingleType) || (_type == TeamCtxInfraTestBlockType) ||
|
||||
(_type == TeamCtxInfraTestOddEvenType) ||
|
||||
(_type == TeamAllToAllTestType) || (_type == TeamFCollectTestType) ||
|
||||
(_type == PingPongTestType) || (_type == BarrierAllTestType) ||
|
||||
(_type == WAVEBarrierAllTestType) || (_type == WGBarrierAllTestType) ||
|
||||
(_type == TeamSyncTestType) || (_type == TeamWAVESyncTestType) ||
|
||||
(_type == TeamWGSyncTestType) || (_type == SyncAllTestType) ||
|
||||
(_type == WAVESyncAllTestType) || (_type == WGSyncAllTestType) ||
|
||||
(_type == RandomAccessTestType) || (_type == PingAllTestType) ||
|
||||
(_type == TeamBarrierTestType) || (_type == TeamWAVEBarrierTestType) ||
|
||||
(_type == TeamWGBarrierTestType) ||
|
||||
(_type == TeamAlltoallmemOnStreamTestType);
|
||||
switch (_type) {
|
||||
case TeamReductionTestType:
|
||||
case TeamBroadcastTestType:
|
||||
case TeamCtxInfraTestType:
|
||||
case TeamCtxInfraTestSingleType:
|
||||
case TeamCtxInfraTestBlockType:
|
||||
case TeamCtxInfraTestOddEvenType:
|
||||
case TeamAllToAllTestType:
|
||||
case TeamFCollectTestType:
|
||||
case PingPongTestType:
|
||||
case BarrierAllTestType:
|
||||
case WAVEBarrierAllTestType:
|
||||
case WGBarrierAllTestType:
|
||||
case TeamSyncTestType:
|
||||
case TeamWAVESyncTestType:
|
||||
case TeamWGSyncTestType:
|
||||
case SyncAllTestType:
|
||||
case WAVESyncAllTestType:
|
||||
case WGSyncAllTestType:
|
||||
case RandomAccessTestType:
|
||||
case PingAllTestType:
|
||||
case TeamBarrierTestType:
|
||||
case TeamWAVEBarrierTestType:
|
||||
case TeamWGBarrierTestType:
|
||||
case TeamAlltoallmemOnStreamTestType:
|
||||
case BarrierAllOnStreamTestType:
|
||||
case TeamBroadcastmemOnStreamTestType:
|
||||
case GetmemOnStreamTestType:
|
||||
case PutmemOnStreamTestType:
|
||||
case PutmemSignalOnStreamTestType:
|
||||
case SignalWaitUntilOnStreamTestType:
|
||||
is_launcher = true;
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
return is_launcher;
|
||||
}
|
||||
|
||||
@@ -114,6 +114,12 @@ enum TestType {
|
||||
TeamCtxInfraTestBlockType = 74,
|
||||
TeamCtxInfraTestOddEvenType = 75,
|
||||
TeamAlltoallmemOnStreamTestType = 76,
|
||||
BarrierAllOnStreamTestType = 77,
|
||||
TeamBroadcastmemOnStreamTestType = 78,
|
||||
GetmemOnStreamTestType = 79,
|
||||
PutmemOnStreamTestType = 80,
|
||||
PutmemSignalOnStreamTestType = 81,
|
||||
SignalWaitUntilOnStreamTestType = 82,
|
||||
};
|
||||
|
||||
enum OpType { PutType = 0, GetType = 1 };
|
||||
|
||||
@@ -112,10 +112,12 @@ TesterArguments::TesterArguments(int argc, char *argv[]) {
|
||||
case TeamBarrierTestType:
|
||||
case TeamWAVEBarrierTestType:
|
||||
case TeamWGBarrierTestType:
|
||||
case BarrierAllOnStreamTestType:
|
||||
case SyncAllTestType:
|
||||
case WAVESyncAllTestType:
|
||||
case WGSyncAllTestType:
|
||||
case TeamSyncTestType:
|
||||
case SignalWaitUntilOnStreamTestType:
|
||||
min_msg_size = 8;
|
||||
max_msg_size = 8;
|
||||
break;
|
||||
@@ -125,6 +127,8 @@ TesterArguments::TesterArguments(int argc, char *argv[]) {
|
||||
max_msg_size = 4;
|
||||
break;
|
||||
case RandomAccessTestType:
|
||||
case TeamAlltoallmemOnStreamTestType:
|
||||
case TeamBroadcastmemOnStreamTestType:
|
||||
min_msg_size = 4;
|
||||
break;
|
||||
case TeamFCollectTestType:
|
||||
@@ -173,23 +177,49 @@ void TesterArguments::get_arguments() {
|
||||
myid = rocshmem_my_pe();
|
||||
|
||||
TestType type = (TestType)algorithm;
|
||||
if ((type != BarrierAllTestType) && (type != WAVEBarrierAllTestType) &&
|
||||
(type != WGBarrierAllTestType) && (type != SyncAllTestType) &&
|
||||
(type != WAVESyncAllTestType) && (type != WGSyncAllTestType) &&
|
||||
(type != TeamSyncTestType) && (type != TeamWAVESyncTestType) &&
|
||||
(type != TeamWGSyncTestType) && (type != TeamAllToAllTestType) &&
|
||||
(type != TeamFCollectTestType) && (type != TeamReductionTestType) &&
|
||||
(type != TeamBroadcastTestType) && (type != PingAllTestType) &&
|
||||
(type != TeamBarrierTestType) && (type != TeamWAVEBarrierTestType) &&
|
||||
(type != TeamWGBarrierTestType) && (type != TeamCtxInfraTestBlockType) &&
|
||||
(type != TeamCtxInfraTestOddEvenType) &&
|
||||
(type != TeamAlltoallmemOnStreamTestType)) {
|
||||
if (numprocs != 2) {
|
||||
if (myid == 0) {
|
||||
std::cerr << "This test requires exactly two processes, we have "
|
||||
<< numprocs << "\n";
|
||||
}
|
||||
exit(-1);
|
||||
// Check if test requires exactly 2 PEs
|
||||
// Tests that support arbitrary number of PEs are excluded
|
||||
bool requires_two_pes = true;
|
||||
switch (type) {
|
||||
// Collective/barrier tests - support any number of PEs
|
||||
case BarrierAllTestType:
|
||||
case WAVEBarrierAllTestType:
|
||||
case WGBarrierAllTestType:
|
||||
case SyncAllTestType:
|
||||
case WAVESyncAllTestType:
|
||||
case WGSyncAllTestType:
|
||||
case TeamSyncTestType:
|
||||
case TeamWAVESyncTestType:
|
||||
case TeamWGSyncTestType:
|
||||
case TeamAllToAllTestType:
|
||||
case TeamFCollectTestType:
|
||||
case TeamReductionTestType:
|
||||
case TeamBroadcastTestType:
|
||||
case PingAllTestType:
|
||||
case TeamBarrierTestType:
|
||||
case TeamWAVEBarrierTestType:
|
||||
case TeamWGBarrierTestType:
|
||||
case TeamCtxInfraTestBlockType:
|
||||
case TeamCtxInfraTestOddEvenType:
|
||||
// On-stream tests - support any number of PEs
|
||||
case TeamAlltoallmemOnStreamTestType:
|
||||
case BarrierAllOnStreamTestType:
|
||||
case TeamBroadcastmemOnStreamTestType:
|
||||
case GetmemOnStreamTestType:
|
||||
case PutmemOnStreamTestType:
|
||||
case PutmemSignalOnStreamTestType:
|
||||
case SignalWaitUntilOnStreamTestType:
|
||||
requires_two_pes = false;
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
if (requires_two_pes && numprocs != 2) {
|
||||
if (myid == 0) {
|
||||
std::cerr << "This test requires exactly two processes, we have "
|
||||
<< numprocs << "\n";
|
||||
}
|
||||
exit(-1);
|
||||
}
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user