Skip to content

test(rocsparse): check wfreduce results only in the specified lanes on gfx8/gfx9 #13388

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
kliegeois wants to merge 1 commit into from
Closed

test(rocsparse): check wfreduce results only in the specified lanes on gfx8/gfx9 #13388

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
7 changes: 5 additions & 2 deletions projects/rocsparse/clients/unittests/unit_test_internal_collective_extras_wfreduce.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,7 @@ using rocsparse_ut::device_vector;
using rocsparse_ut::launch_single_warp;
using rocsparse_ut::launch_warp_by_size;
using rocsparse_ut::to_host;
using rocsparse_ut::wfreduce_first_valid_lane;

using namespace rocsparse_ut_collective_extras;

Expand Down Expand Up @@ -104,7 +105,9 @@ namespace
}

// Host reference mirrors the exact xor-butterfly (stride wf/2 down to SUB) at
// the device's runtime wavefront width; every lane's result is checked.
// the device's runtime wavefront width; every lane's result is checked on the
// shuffle path, the last SUB lanes on the gfx8/gfx9 DPP path (see
// rocsparse_ut::wfreduce_first_valid_lane).
template <uint32_t SUB, typename T, typename Gen>
void run_wfreduce_partial_sum(Gen gen)
{
Expand All @@ -130,7 +133,7 @@ namespace
d_out.ptr),
hipSuccess);
auto h = to_host(d_out);
for(uint32_t l = 0; l < wf; ++l)
for(uint32_t l = wfreduce_first_valid_lane(wf, SUB); l < wf; ++l)
EXPECT_DOUBLE_EQ(h[l], cur[l]);
}
} // namespace
Expand Down
34 changes: 19 additions & 15 deletions projects/rocsparse/clients/unittests/unit_test_internal_collectives_warpreduce.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,6 +49,7 @@
using rocsparse_ut::device_vector;
using rocsparse_ut::launch_warp_by_size;
using rocsparse_ut::to_host;
using rocsparse_ut::wfreduce_first_valid_lane;

using namespace rocsparse_ut_collectives;

Expand Down Expand Up @@ -83,10 +84,11 @@ namespace
const int lane = threadIdx.x;
out[lane] = rocsparse::wfreduce_partial_sum<WFSIZE, SUB>(in[lane]);
}
// wfreduce_sum<WFSIZE>(x): all-reduce sum across one wavefront; every lane
// returns the total. Runs on the device's own wavefront width (32 or 64) by
// dispatching to the matching instantiation; the host reference is the exact
// sum of the wf lane values produced by `gen`.
// wfreduce_sum<WFSIZE>(x): sum across one wavefront; the total is returned in
// every lane (in lane wf-1 only on the gfx8/gfx9 DPP path). Runs on the
// device's own wavefront width (32 or 64) by dispatching to the matching
// instantiation; the host reference is the exact sum of the wf lane values
// produced by `gen`.
template <typename T, typename Gen>
void run_wfreduce_sum(Gen gen)
{
Expand All @@ -104,12 +106,12 @@ namespace
launch_warp_by_size(k_wfreduce_sum<32, T>, k_wfreduce_sum<64, T>, d_in.ptr, d_out.ptr),
hipSuccess);
auto h = to_host(d_out);
for(uint32_t l = 0; l < wf; ++l)
for(uint32_t l = wfreduce_first_valid_lane(wf, 1); l < wf; ++l)
expect_close(h[l], ref);
}
// wfreduce_max<WFSIZE>(&v): all-reduce max across one wavefront; every lane
// ends holding the maximum. Host reference is std::max_element over the wf
// lane values from `gen`.
// wfreduce_max<WFSIZE>(&v): max across one wavefront; every lane ends holding
// the maximum (lane wf-1 only on the gfx8/gfx9 DPP path). Host reference is
// std::max_element over the wf lane values from `gen`.
template <typename T, typename Gen>
void run_wfreduce_max(Gen gen)
{
Expand All @@ -125,10 +127,10 @@ namespace
launch_warp_by_size(k_wfreduce_max<32, T>, k_wfreduce_max<64, T>, d_in.ptr, d_out.ptr),
hipSuccess);
auto h = to_host(d_out);
for(uint32_t l = 0; l < wf; ++l)
for(uint32_t l = wfreduce_first_valid_lane(wf, 1); l < wf; ++l)
expect_close(h[l], ref);
}
// wfreduce_min<WFSIZE>(&v): all-reduce min across one wavefront.
// wfreduce_min<WFSIZE>(&v): min across one wavefront (same lane contract as max).
template <typename T, typename Gen>
void run_wfreduce_min(Gen gen)
{
Expand All @@ -144,7 +146,7 @@ namespace
launch_warp_by_size(k_wfreduce_min<32, T>, k_wfreduce_min<64, T>, d_in.ptr, d_out.ptr),
hipSuccess);
auto h = to_host(d_out);
for(uint32_t l = 0; l < wf; ++l)
for(uint32_t l = wfreduce_first_valid_lane(wf, 1); l < wf; ++l)
expect_close(h[l], ref);
}
// Distinct, non-monotone exact values across the wavefront so max/min land on
Expand All @@ -165,7 +167,8 @@ namespace
}
// wfreduce_partial_sum<WFSIZE, SUB>(x): xor-butterfly that sums within each
// SUB-lane sub-group. Host reference mirrors the exact butterfly (halving the
// stride from wf/2 down to SUB), so every lane is checked, not just one.
// stride from wf/2 down to SUB); every lane is checked on the shuffle path and
// the last SUB lanes on the gfx8/gfx9 DPP path (see wfreduce_first_valid_lane).
template <uint32_t SUB, typename T, typename Gen>
void run_wfreduce_partial_sum(Gen gen)
{
Expand All @@ -191,13 +194,14 @@ namespace
d_out.ptr),
hipSuccess);
auto h = to_host(d_out);
for(uint32_t l = 0; l < wf; ++l)
for(uint32_t l = wfreduce_first_valid_lane(wf, SUB); l < wf; ++l)
expect_close(h[l], cur[l]);
}
} // namespace

// Input: exact small integers per lane. Expected: every lane == sum over the
// wavefront. One case per element type the routine supports.
// Input: exact small integers per lane. Expected: every lane (the last lane on
// gfx8/gfx9) == sum over the wavefront. One case per element type the routine
// supports.
TEST(internal_collectives_wfreduce_sum, i32)
{
run_wfreduce_sum<int32_t>([](uint32_t l) { return static_cast<int32_t>((l % 7) + 1); });
Expand Down
29 changes: 29 additions & 0 deletions projects/rocsparse/clients/unittests/unit_test_utils.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,8 @@

#include "rocsparse.h"

#include <cstdint>
#include <cstring>
#include <gtest/gtest.h>
#include <hip/hip_runtime.h>
#include <vector>
Expand Down Expand Up @@ -222,6 +224,33 @@ namespace rocsparse_ut
return props.warpSize;
}

// True on gfx8 / gfx9, where the rocsparse::wfreduce_* routines are built on
// DPP row_shr / row_bcast steps (ROCSPARSE_USE_MOVE_DPP). Those leave the
// result only in the TOP lanes of the wavefront: lane wf-1 for
// wfreduce_sum / _max / _min, and the last SUB lanes for
// wfreduce_partial_sum<WFSIZE, SUB> (the lanes callers read). On every other
// architecture the shuffle-based variants return the result in every lane.
// ROCSPARSE_USE_MOVE_DPP only exists in the device pass, so the host tests
// select the architecture from the device name instead.
inline bool device_uses_dpp_wfreduce()
{
int dev = 0;
hipDeviceProp_t props = {};
if(hipGetDevice(&dev) != hipSuccess || hipGetDeviceProperties(&props, dev) != hipSuccess)
return false;
return std::strncmp(props.gcnArchName, "gfx8", 4) == 0
|| std::strncmp(props.gcnArchName, "gfx9", 4) == 0;
}

// First lane of a wavefront of `wf` lanes whose wfreduce_* result is
// specified for a reduction leaving `sub` result lanes (sub == 1 for
// sum / max / min): lane 0 (every lane) on the shuffle path, lane wf-sub on
// the DPP path. Check lanes [first, wf).
inline uint32_t wfreduce_first_valid_lane(uint32_t wf, uint32_t sub)
{
return device_uses_dpp_wfreduce() ? wf - sub : 0;
}

#ifdef __HIPCC__
// ------------------------------------------------------------------------
// Single-block / single-warp launch helpers
Expand Down
Loading