AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

Convert simple fine tests into parameterized tests

Przeglądaj źródła 
[ROCm/rocshmem commit: 749d9f0781]
This commit is contained in:
Brandon Potter
2024-10-09 12:07:37 -05:00
rodzic e21e9a8fb6
commit 74dec9374d
4 zmienionych plików z 213 dodań i 1058 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.cpp
+14 -14
Wyświetl plik
@@ -25,15 +25,15 @@
using namespace rocshmem;
TEST_P(Degenerate, ptr_check) {
TEST_P(DegenerateSimpleCoarse, ptr_check) {
ASSERT_NE(heap_mem_.get_ptr(), nullptr);
}
TEST_P(Degenerate, MPI_num_pes) {
TEST_P(DegenerateSimpleCoarse, MPI_num_pes) {
ASSERT_EQ(mpi_.num_pes(), 2);
}
TEST_P(Degenerate, IPC_bases) {
TEST_P(DegenerateSimpleCoarse, IPC_bases) {
ASSERT_EQ(mpi_.num_pes(), 2);
ASSERT_NE(ipc_impl_.ipc_bases, nullptr);
for(int i{0}; i < mpi_.num_pes(); i++) {
@@ -41,28 +41,28 @@ TEST_P(Degenerate, IPC_bases) {
}
}
TEST_P(Degenerate, golden_1048576_int) {
TEST_P(DegenerateSimpleCoarse, golden_1048576_int) {
iota_golden(1048576);
validate_golden(1048576);
}
INSTANTIATE_TEST_SUITE_P(
IPCImplSimpleCoarseTestFixture,
Degenerate,
DegenerateSimpleCoarse,
::testing::Values(
std::make_tuple(1, 1, 1))
);
//=============================================================================
TEST_P(ParameterizedBlock, write) {
TEST_P(ParameterizedBlockSimpleCoarse, write) {
dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
size_t size = std::get<2>(GetParam());
write(grid, block, size);
}
TEST_P(ParameterizedBlock, read) {
TEST_P(ParameterizedBlockSimpleCoarse, read) {
dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
size_t size = std::get<2>(GetParam());
@@ -71,7 +71,7 @@ TEST_P(ParameterizedBlock, read) {
INSTANTIATE_TEST_SUITE_P(
IPCImplSimpleCoarseTestFixture,
ParameterizedBlock,
ParameterizedBlockSimpleCoarse,
::testing::Values(
std::make_tuple(1, 1024, 32), // 0
std::make_tuple(1, 1, 1048576), // 1
@@ -90,14 +90,14 @@ INSTANTIATE_TEST_SUITE_P(
//=============================================================================
TEST_P(ParameterizedWarp, write) {
TEST_P(ParameterizedWarpSimpleCoarse, write) {
dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
size_t size = std::get<2>(GetParam());
write(grid, block, size);
}
TEST_P(ParameterizedWarp, read) {
TEST_P(ParameterizedWarpSimpleCoarse, read) {
dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
size_t size = std::get<2>(GetParam());
@@ -106,7 +106,7 @@ TEST_P(ParameterizedWarp, read) {
INSTANTIATE_TEST_SUITE_P(
IPCImplSimpleCoarseTestFixture,
ParameterizedWarp,
ParameterizedWarpSimpleCoarse,
::testing::Values(
std::make_tuple(1, 64, 1), // 0
std::make_tuple(1, 64, 32), // 1
@@ -179,14 +179,14 @@ INSTANTIATE_TEST_SUITE_P(
//=============================================================================
TEST_P(ParameterizedThread, write) {
TEST_P(ParameterizedThreadSimpleCoarse, write) {
dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
size_t size = std::get<2>(GetParam());
write(grid, block, size);
}
TEST_P(ParameterizedThread, read) {
TEST_P(ParameterizedThreadSimpleCoarse, read) {
dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
size_t size = std::get<2>(GetParam());
@@ -195,7 +195,7 @@ TEST_P(ParameterizedThread, read) {
INSTANTIATE_TEST_SUITE_P(
IPCImplSimpleCoarseTestFixture,
ParameterizedThread,
ParameterizedThreadSimpleCoarse,
::testing::Values(
std::make_tuple(1, 1, 1048576))
);
projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp
+8 -8
Wyświetl plik
@@ -186,32 +186,32 @@ class IPCImplSimpleCoarse : public ::testing::TestWithParam<std::tuple<int, int,
HIPAllocator hip_allocator_ {};
};
class Degenerate : public IPCImplSimpleCoarse {
class DegenerateSimpleCoarse : public IPCImplSimpleCoarse {
public:
~Degenerate() override {};
~DegenerateSimpleCoarse() override {};
};
class ParameterizedBlock : public IPCImplSimpleCoarse {
class ParameterizedBlockSimpleCoarse : public IPCImplSimpleCoarse {
public:
~ParameterizedBlock() override {};
~ParameterizedBlockSimpleCoarse() override {};
void copy(IPCImplSimpleCoarse::TestType test, dim3 grid, dim3 block) override {
execute(test, kernel_simple_coarse_copy_block, grid, block);
}
};
class ParameterizedWarp : public IPCImplSimpleCoarse {
class ParameterizedWarpSimpleCoarse : public IPCImplSimpleCoarse {
public:
~ParameterizedWarp() override {};
~ParameterizedWarpSimpleCoarse() override {};
void copy(IPCImplSimpleCoarse::TestType test, dim3 grid, dim3 block) override {
execute(test, kernel_simple_coarse_copy_warp, grid, block);
}
};
class ParameterizedThread : public IPCImplSimpleCoarse {
class ParameterizedThreadSimpleCoarse : public IPCImplSimpleCoarse {
public:
~ParameterizedThread() override {};
~ParameterizedThreadSimpleCoarse() override {};
void copy(IPCImplSimpleCoarse::TestType test, dim3 grid, dim3 block) override {
execute(test, kernel_simple_coarse_copy, grid, block);
projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.cpp
+148 -985
Wyświetl plik
Plik diff jest za duży Load Diff
projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp
+43 -51
Wyświetl plik
@@ -36,13 +36,13 @@
namespace rocshmem {
const uint32_t SIGNAL_OFFSET {67108864};
enum TestType {
READ = 0,
WRITE = 1
};
const uint32_t SIGNAL_OFFSET {67108864};
__device__
void
validator(bool *error, int *golden, int *dest, size_t bytes) {
@@ -89,7 +89,7 @@ kernel_simple_fine_copy(IpcImpl *ipc_impl, bool *error, int *golden, int *src, i
template <typename NotifierT>
__global__
void
kernel_simple_fine_copy_wg(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) {
kernel_simple_fine_copy_block(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) {
if (!blockIdx.x) {
ipc_impl->ipcCopy_wg(dest, src, bytes);
ipc_impl->ipcFence();
@@ -108,7 +108,7 @@ kernel_simple_fine_copy_wg(IpcImpl *ipc_impl, bool *error, int *golden, int *src
template <typename NotifierT>
__global__
void
kernel_simple_fine_copy_wave(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) {
kernel_simple_fine_copy_warp(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) {
if (!blockIdx.x && threadIdx.x < 64) {
ipc_impl->ipcCopy_wave(dest, src, bytes);
ipc_impl->ipcFence();
@@ -125,7 +125,7 @@ kernel_simple_fine_copy_wave(IpcImpl *ipc_impl, bool *error, int *golden, int *s
}
}
class IPCImplSimpleFineTestFixture : public ::testing::Test {
class IPCImplSimpleFine : public ::testing::TestWithParam<std::tuple<int, int, int>> {
using HEAP_T = HeapMemory<HIPDefaultFinegrainedAllocator>;
using MPI_T = RemoteHeapInfo<CommunicatorMPI>;
using NotifierT = Notifier<detail::atomic::memory_scope_agent>;
@@ -134,7 +134,7 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test {
using FN_T2 = void (*)(bool*, int*, int*, size_t, NotifierT*);
public:
IPCImplSimpleFineTestFixture() {
IPCImplSimpleFine() {
ipc_impl_.ipcHostInit(mpi_.my_pe(), mpi_.get_heap_bases() , MPI_COMM_WORLD);
assert(ipc_impl_dptr_ == nullptr);
@@ -146,7 +146,7 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test {
*error_dptr_ = false;
}
~IPCImplSimpleFineTestFixture() {
~IPCImplSimpleFine() {
if (ipc_impl_dptr_) {
hip_allocator_.deallocate(ipc_impl_dptr_);
}
@@ -170,6 +170,10 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test {
CHECK_HIP(hipStreamSynchronize(nullptr));
}
virtual void copy(TestType test, dim3 grid, dim3 block) {
FAIL();
}
void write(const dim3 grid, const dim3 block, size_t elems) {
iota_golden(elems);
initialize_signal(WRITE);
@@ -178,22 +182,6 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test {
check_device_validation_errors(WRITE);
}
void write_wg(const dim3 grid, const dim3 block, size_t elems) {
iota_golden(elems);
initialize_signal(WRITE);
initialize_src_buffer(WRITE);
copy_wg(WRITE, grid, block);
check_device_validation_errors(WRITE);
}
void write_wave(const dim3 grid, const dim3 block, size_t elems) {
iota_golden(elems);
initialize_signal(WRITE);
initialize_src_buffer(WRITE);
copy_wave(WRITE, grid, block);
check_device_validation_errors(WRITE);
}
void read(const dim3 grid, const dim3 block, size_t elems) {
iota_golden(elems);
initialize_signal(READ);
@@ -202,22 +190,6 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test {
check_device_validation_errors(READ);
}
void read_wg(const dim3 grid, const dim3 block, size_t elems) {
iota_golden(elems);
initialize_signal(READ);
initialize_src_buffer(READ);
copy_wg(READ, grid, block);
check_device_validation_errors(READ);
}
void read_wave(const dim3 grid, const dim3 block, size_t elems) {
iota_golden(elems);
initialize_signal(READ);
initialize_src_buffer(READ);
copy_wave(READ, grid, block);
check_device_validation_errors(READ);
}
void iota_golden(size_t elems) {
golden_.resize(elems);
std::iota(golden_.begin(), golden_.end(), 0);
@@ -285,18 +257,6 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test {
mpi_.barrier();
}
void copy(TestType test, dim3 grid, dim3 block) {
execute(test, kernel_simple_fine_copy, grid, block);
}
void copy_wg(TestType test, dim3 grid, dim3 block) {
execute(test, kernel_simple_fine_copy_wg, grid, block);
}
void copy_wave(TestType test, dim3 grid, dim3 block) {
execute(test, kernel_simple_fine_copy_wave, grid, block);
}
void check_device_validation_errors(TestType test) {
if (!pe_validates_dest_buffer(test)) {
return;
@@ -339,6 +299,38 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test {
bool *error_dptr_ {nullptr};
};
class DegenerateSimpleFine : public IPCImplSimpleFine {
public:
~DegenerateSimpleFine() override {};
};
class ParameterizedBlockSimpleFine : public IPCImplSimpleFine {
public:
~ParameterizedBlockSimpleFine() override {};
void copy(TestType test, dim3 grid, dim3 block) override {
execute(test, kernel_simple_fine_copy_block, grid, block);
}
};
class ParameterizedWarpSimpleFine : public IPCImplSimpleFine {
public:
~ParameterizedWarpSimpleFine() override {};
void copy(TestType test, dim3 grid, dim3 block) override {
execute(test, kernel_simple_fine_copy_warp, grid, block);
}
};
class ParameterizedThreadSimpleFine : public IPCImplSimpleFine {
public:
~ParameterizedThreadSimpleFine() override {};
void copy(TestType test, dim3 grid, dim3 block) override {
execute(test, kernel_simple_fine_copy, grid, block);
}
};
} // namespace rocshmem
#endif // ROCSHMEM_IPC_IMPL_SIMPLE_FINE_GTEST_HPP