mirror of
https://github.com/microsoft/mscclpp.git
synced 2026-05-13 09:46:00 +00:00
93f6eeaa6b
- Removes the GTest dependency, replacing it with a minimal custom framework (`test/framework.*`) that covers only what the tests actually use — a unified `TEST()` macro with SFINAE-based fixture auto-detection, `EXPECT_*`/`ASSERT_*` assertions, environments, and setup/teardown. - `--exclude-perf-tests` flag and substring-based negative filtering - `MSCCLPP_ENABLE_COVERAGE` CMake option with gcov/lcov; CI uploads to Codecov - Merges standalone `test/perf/` into main test targets - Refactors Azure pipelines to reduce redundancies & make more readable --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: Changho Hwang <changhohwang@microsoft.com>
76 lines
2.0 KiB
Plaintext
76 lines
2.0 KiB
Plaintext
// Copyright (c) Microsoft Corporation.
|
|
// Licensed under the MIT License.
|
|
|
|
#include <mscclpp/fifo.hpp>
|
|
#include <mscclpp/gpu_utils.hpp>
|
|
#include <mscclpp/numa.hpp>
|
|
#include <mscclpp/utils.hpp>
|
|
|
|
#include "../framework.hpp"
|
|
#include "utils_internal.hpp"
|
|
|
|
#define ITER 10000 // should be larger than the FIFO size for proper testing
|
|
|
|
__constant__ mscclpp::FifoDeviceHandle gFifoTestFifoDeviceHandle;
|
|
__global__ void kernelFifoTest() {
|
|
if (threadIdx.x + blockIdx.x * blockDim.x != 0) return;
|
|
|
|
mscclpp::FifoDeviceHandle& fifo = gFifoTestFifoDeviceHandle;
|
|
mscclpp::ProxyTrigger trigger;
|
|
for (uint64_t i = 1; i < ITER + 1; ++i) {
|
|
trigger.fst = i;
|
|
trigger.snd = i;
|
|
uint64_t curFifoHead = fifo.push(trigger);
|
|
if (i % fifo.size == 0) {
|
|
fifo.sync(curFifoHead);
|
|
}
|
|
}
|
|
}
|
|
|
|
TEST(FifoTest, Fifo) {
|
|
int cudaNum;
|
|
MSCCLPP_CUDATHROW(cudaGetDevice(&cudaNum));
|
|
int numaNode = mscclpp::getDeviceNumaNode(cudaNum);
|
|
mscclpp::numaBind(numaNode);
|
|
|
|
mscclpp::Fifo hostFifo;
|
|
if (hostFifo.size() >= ITER) {
|
|
FAIL() << "ITER is too small for proper testing.";
|
|
}
|
|
|
|
mscclpp::FifoDeviceHandle devFifo = hostFifo.deviceHandle();
|
|
MSCCLPP_CUDATHROW(cudaMemcpyToSymbol(gFifoTestFifoDeviceHandle, &devFifo, sizeof(devFifo)));
|
|
|
|
kernelFifoTest<<<1, 1>>>();
|
|
MSCCLPP_CUDATHROW(cudaGetLastError());
|
|
|
|
mscclpp::ProxyTrigger trigger;
|
|
trigger.fst = 0;
|
|
trigger.snd = 0;
|
|
|
|
uint64_t spin = 0;
|
|
mscclpp::Timer timer(3);
|
|
for (uint64_t i = 0; i < ITER; ++i) {
|
|
trigger = hostFifo.poll();
|
|
while (trigger.fst == 0 || trigger.snd == 0) {
|
|
trigger = hostFifo.poll();
|
|
|
|
if (spin++ > 1000000) {
|
|
FAIL() << "Polling is stuck.";
|
|
}
|
|
}
|
|
// see `src/proxy.cc` for the reason of this line
|
|
trigger.snd ^= ((uint64_t)1 << (uint64_t)63);
|
|
ASSERT_TRUE(trigger.fst == (i + 1));
|
|
ASSERT_TRUE(trigger.snd == (i + 1));
|
|
hostFifo.pop();
|
|
spin = 0;
|
|
}
|
|
|
|
std::stringstream ss;
|
|
ss << "FifoTest.Fifo: " << (float)timer.elapsed() / ITER << " us/iter\n";
|
|
std::cout << ss.str();
|
|
|
|
MSCCLPP_CUDATHROW(cudaDeviceSynchronize());
|
|
}
|