diff --git a/projects/rocsparse/clients/unittests/unit_test_internal_collective_extras_wfreduce.cpp b/projects/rocsparse/clients/unittests/unit_test_internal_collective_extras_wfreduce.cpp index a62428b51974..3eb0d7f179a6 100644 --- a/projects/rocsparse/clients/unittests/unit_test_internal_collective_extras_wfreduce.cpp +++ b/projects/rocsparse/clients/unittests/unit_test_internal_collective_extras_wfreduce.cpp @@ -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; @@ -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 void run_wfreduce_partial_sum(Gen gen) { @@ -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 diff --git a/projects/rocsparse/clients/unittests/unit_test_internal_collectives_warpreduce.cpp b/projects/rocsparse/clients/unittests/unit_test_internal_collectives_warpreduce.cpp index 7b3c1f1e8f9c..d912d6d79148 100644 --- a/projects/rocsparse/clients/unittests/unit_test_internal_collectives_warpreduce.cpp +++ b/projects/rocsparse/clients/unittests/unit_test_internal_collectives_warpreduce.cpp @@ -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; @@ -83,10 +84,11 @@ namespace const int lane = threadIdx.x; out[lane] = rocsparse::wfreduce_partial_sum(in[lane]); } - // wfreduce_sum(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(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 void run_wfreduce_sum(Gen gen) { @@ -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(&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(&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 void run_wfreduce_max(Gen gen) { @@ -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(&v): all-reduce min across one wavefront. + // wfreduce_min(&v): min across one wavefront (same lane contract as max). template void run_wfreduce_min(Gen gen) { @@ -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 @@ -165,7 +167,8 @@ namespace } // wfreduce_partial_sum(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 void run_wfreduce_partial_sum(Gen gen) { @@ -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([](uint32_t l) { return static_cast((l % 7) + 1); }); diff --git a/projects/rocsparse/clients/unittests/unit_test_utils.hpp b/projects/rocsparse/clients/unittests/unit_test_utils.hpp index 8c89fbdd7603..8a4c7a400a2c 100644 --- a/projects/rocsparse/clients/unittests/unit_test_utils.hpp +++ b/projects/rocsparse/clients/unittests/unit_test_utils.hpp @@ -34,6 +34,8 @@ #include "rocsparse.h" +#include +#include #include #include #include @@ -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 (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