AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge pull request #54 from Yiltan/sig-ops

Browse Source 
Added SHMEM Signalling Operations
This commit is contained in:
Yiltan
2024-11-26 15:38:12 -05:00
committed by GitHub
parent 46f0b42ac3 f710a301fe
commit 0c5524d7df
21 changed files with 865 additions and 13 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/functional_tests/CMakeLists.txt
+2
View File
@@ -53,6 +53,8 @@ target_sources(
swarm_tester.cpp
random_access_tester.cpp
shmem_ptr_tester.cpp
signaling_operations_tester.cpp
signaling_operations_tester.hpp
extended_primitives.cpp
empty_tester.cpp
wave_level_primitives.cpp
tests/functional_tests/signaling_operations_tester.cpp
+163
View File
@@ -0,0 +1,163 @@
/******************************************************************************
* Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to
* deal in the Software without restriction, including without limitation the
* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
* sell copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in
* all copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
* FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
* IN THE SOFTWARE.
*****************************************************************************/
#include "signaling_operations_tester.hpp"
#include <roc_shmem/roc_shmem.hpp>
using namespace rocshmem;
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void SignalingOperationsTest(int loop, int skip, uint64_t *timer, char *s_buf,
char *r_buf, int size, uint64_t *sig_addr,
uint64_t *fetched_value,
TestType type, ShmemContextType ctx_type) {
__shared__ roc_shmem_ctx_t ctx;
roc_shmem_wg_init();
roc_shmem_wg_ctx_create(ctx_type, &ctx);
uint64_t start;
uint64_t signal = 0;
int sig_op = ROC_SHMEM_SIGNAL_SET;
for (int i = 0; i < loop + skip; i++) {
if (i == skip) {
__syncthreads();
start = roc_shmem_timer();
}
switch (type) {
case PutSignalTestType:
roc_shmem_ctx_putmem_signal(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
break;
case WGPutSignalTestType:
roc_shmem_ctx_putmem_signal_wg(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
break;
case WAVEPutSignalTestType:
roc_shmem_ctx_putmem_signal_wave(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
break;
case PutSignalNBITestType:
roc_shmem_ctx_putmem_signal_nbi(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
break;
case WGPutSignalNBITestType:
roc_shmem_ctx_putmem_signal_nbi_wg(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
break;
case WAVEPutSignalNBITestType:
roc_shmem_ctx_putmem_signal_nbi_wave(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1);
break;
case SignalFetchTestType:
*fetched_value = roc_shmem_signal_fetch(sig_addr);
break;
case WGSignalFetchTestType:
*fetched_value = roc_shmem_signal_fetch_wg(sig_addr);
break;
case WAVESignalFetchTestType:
*fetched_value = roc_shmem_signal_fetch_wave(sig_addr);
break;
default:
break;
}
}
roc_shmem_ctx_quiet(ctx);
__syncthreads();
if (hipThreadIdx_x == 0) {
timer[hipBlockIdx_x] = roc_shmem_timer() - start;
}
roc_shmem_wg_ctx_destroy(&ctx);
roc_shmem_wg_finalize();
}
/******************************************************************************
* HOST TESTER CLASS METHODS
*****************************************************************************/
SignalingOperationsTester::SignalingOperationsTester(TesterArguments args) : Tester(args) {
s_buf = (char *)roc_shmem_malloc(args.max_msg_size * args.wg_size);
r_buf = (char *)roc_shmem_malloc(args.max_msg_size * args.wg_size);
sig_addr = (uint64_t *)roc_shmem_malloc(sizeof(uint64_t));
CHECK_HIP(hipMallocManaged(&fetched_value, sizeof(uint64_t), hipMemAttachHost));
}
SignalingOperationsTester::~SignalingOperationsTester() {
roc_shmem_free(s_buf);
roc_shmem_free(r_buf);
roc_shmem_free(sig_addr);
CHECK_HIP(hipFree(fetched_value));
}
void SignalingOperationsTester::resetBuffers(uint64_t size) {
memset(s_buf, '0', args.max_msg_size * args.wg_size);
memset(r_buf, '1', args.max_msg_size * args.wg_size);
*fetched_value = -1;
*sig_addr = args.myid + 123;
}
void SignalingOperationsTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
uint64_t size) {
size_t shared_bytes = 0;
hipLaunchKernelGGL(SignalingOperationsTest, gridSize, blockSize, shared_bytes, stream,
loop, args.skip, timer, s_buf, r_buf, size, sig_addr, fetched_value,
_type, _shmem_context);
num_msgs = (loop + args.skip) * gridSize.x;
num_timed_msgs = loop;
}
void SignalingOperationsTester::verifyResults(uint64_t size) {
int check_data_id = (_type == PutSignalTestType ||
_type == PutSignalNBITestType ||
_type == WAVEPutSignalTestType ||
_type == WAVEPutSignalNBITestType ||
_type == WGPutSignalTestType ||
_type == WGPutSignalNBITestType)
? 1 : -1; // do not check if it doesn't match a test
int check_fetched_value_id = (_type == SignalFetchTestType ||
_type == WAVESignalFetchTestType ||
_type == WGSignalFetchTestType)
? 0 : -1; // do not check if it doesn't match a test
if (args.myid == check_data_id) {
for (int i = 0; i < size; i++) {
if (r_buf[i] != '0') {
fprintf(stderr, "Data validation error at idx %d\n", i);
fprintf(stderr, "Got %c, Expected %c\n", r_buf[i], '0');
exit(-1);
}
}
}
if (args.myid == check_fetched_value_id) {
uint64_t value = *fetched_value;
uint64_t expected_value = (args.myid + 123);
if (value != expected_value) {
fprintf(stderr, "Fetched Value %lu, Expected %lu\n", value, expected_value);
exit(-1);
}
}
}
tests/functional_tests/signaling_operations_tester.hpp
+50
View File
@@ -0,0 +1,50 @@
/******************************************************************************
* Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to
* deal in the Software without restriction, including without limitation the
* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
* sell copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in
* all copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
* FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
* IN THE SOFTWARE.
*****************************************************************************/
#ifndef _SIGNALING_OPERATIONS_
#define _SIGNALING_OPERATIONS_
#include "tester.hpp"
/******************************************************************************
* HOST TESTER CLASS
*****************************************************************************/
class SignalingOperationsTester : public Tester {
public:
explicit SignalingOperationsTester(TesterArguments args);
virtual ~SignalingOperationsTester();
protected:
virtual void resetBuffers(uint64_t size) override;
virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop,
uint64_t size) override;
virtual void verifyResults(uint64_t size) override;
char *s_buf = nullptr;
char *r_buf = nullptr;
uint64_t *sig_addr;
uint64_t *fetched_value;
};
#endif
tests/functional_tests/tester.cpp
+37
View File
@@ -44,6 +44,7 @@
#include "primitive_tester.hpp"
#include "random_access_tester.hpp"
#include "shmem_ptr_tester.hpp"
#include "signaling_operations_tester.hpp"
#include "swarm_tester.hpp"
#include "sync_tester.hpp"
#include "team_broadcast_tester.hpp"
@@ -419,6 +420,42 @@ std::vector<Tester*> Tester::create(TesterArguments args) {
}
testers.push_back(new WaveLevelPrimitiveTester(args));
return testers;
case PutSignalTestType:
if (rank == 0) std::cout << "Putmem Signal ###" << std::endl;
testers.push_back(new SignalingOperationsTester(args));
return testers;
case WGPutSignalTestType:
if (rank == 0) std::cout << "WG Putmem Signal ###" << std::endl;
testers.push_back(new SignalingOperationsTester(args));
return testers;
case WAVEPutSignalTestType:
if (rank == 0) std::cout << "Wave Putmem Signal ###" << std::endl;
testers.push_back(new SignalingOperationsTester(args));
return testers;
case PutSignalNBITestType:
if (rank == 0) std::cout << "Non-Blocking Putmem Signal ###" << std::endl;
testers.push_back(new SignalingOperationsTester(args));
return testers;
case WGPutSignalNBITestType:
if (rank == 0) std::cout << "Non-Blocking WG Putmem Signal ###" << std::endl;
testers.push_back(new SignalingOperationsTester(args));
return testers;
case WAVEPutSignalNBITestType:
if (rank == 0) std::cout << "Non-Blocking Wave Putmem Signal ###" << std::endl;
testers.push_back(new SignalingOperationsTester(args));
return testers;
case SignalFetchTestType:
if (rank == 0) std::cout << "Signal Fetch ###" << std::endl;
testers.push_back(new SignalingOperationsTester(args));
return testers;
case WGSignalFetchTestType:
if (rank == 0) std::cout << "WG Signal Fetch ###" << std::endl;
testers.push_back(new SignalingOperationsTester(args));
return testers;
case WAVESignalFetchTestType:
if (rank == 0) std::cout << "Wave Signal Fetch ###" << std::endl;
testers.push_back(new SignalingOperationsTester(args));
return testers;
default:
if (rank == 0) std::cout << "Unknown ###" << std::endl;
testers.push_back(new PrimitiveTester(args));
tests/functional_tests/tester.hpp
+10 -1
View File
@@ -84,7 +84,16 @@ enum TestType {
AMO_AndTestType = 49,
AMO_OrTestType = 50,
AMO_XorTestType = 51,
PingAllTestType = 52
PingAllTestType = 52,
PutSignalTestType = 53,
WGPutSignalTestType = 54,
WAVEPutSignalTestType = 55,
PutSignalNBITestType = 56,
WGPutSignalNBITestType = 57,
WAVEPutSignalNBITestType = 58,
SignalFetchTestType = 59,
WGSignalFetchTestType = 60,
WAVESignalFetchTestType = 61,
};
enum OpType { PutType = 0, GetType = 1 };