SWDEV-420237 - Add tests for reduce sync operations (#102)
Dieser Commit ist enthalten in:
Hernandez, Gerardo
committet von
GitHub
GitHub
Ursprung
d9cb3c7fb1
Commit
001affc6ae
@@ -30,9 +30,6 @@ int main(int argc, char** argv) {
|
||||
| Opt(cmd_options.progress)
|
||||
["-P"]["--progress"]
|
||||
("Show progress bar when running performance tests")
|
||||
| Opt(cmd_options.cg_extended_run, "cg_extened_run")
|
||||
["-E"]["--cg-extended-run"]
|
||||
("TODO: Description goes here")
|
||||
| Opt(cmd_options.cg_iterations, "cg_iterations")
|
||||
["-C"]["--cg-iterations"]
|
||||
("Number of iterations used for cooperative groups sync tests (default: 5)")
|
||||
@@ -42,6 +39,12 @@ int main(int argc, char** argv) {
|
||||
| Opt(cmd_options.accuracy_max_memory, "accuracy_max_memory")
|
||||
["-M"]["--accuracy-max-memory"]
|
||||
("Percentage of global device memory allowed for math accuracy tests (default: 80%)")
|
||||
| Opt(cmd_options.reduce_iterations, "reduce_iterations")
|
||||
["-R"]["--reduce-iterations"]
|
||||
("Number of iterations for fuzzing reduce operations (default: 1)")
|
||||
| Opt(cmd_options.reduce_input_size, "reduce_input_size")
|
||||
["-Z"]["--reduce-input-size"]
|
||||
("Size of the input for the reduce sync operations performance test (megabytes) (default: 50)")
|
||||
;
|
||||
// clang-format on
|
||||
|
||||
|
||||
@@ -28,11 +28,12 @@ THE SOFTWARE.
|
||||
struct CmdOptions {
|
||||
int iterations = 10;
|
||||
int warmups = 100;
|
||||
int cg_extended_run = 5;
|
||||
int cg_iterations = 5;
|
||||
bool no_display = false;
|
||||
bool progress = false;
|
||||
uint64_t accuracy_iterations = std::numeric_limits<uint32_t>::max() + 1ull;
|
||||
uint64_t reduce_iterations = 1;
|
||||
uint64_t reduce_input_size = 50;
|
||||
int accuracy_max_memory = 80;
|
||||
};
|
||||
|
||||
|
||||
@@ -55,7 +55,8 @@ template <typename T> class LinearAllocGuard {
|
||||
|
||||
LinearAllocGuard(const LinearAllocs allocation_type, const size_t size,
|
||||
const unsigned int flags = 0u)
|
||||
: allocation_type_{allocation_type} {
|
||||
: allocation_type_{allocation_type},
|
||||
size_{size} {
|
||||
switch (allocation_type_) {
|
||||
case LinearAllocs::malloc:
|
||||
ptr_ = host_ptr_ = reinterpret_cast<T*>(malloc(size));
|
||||
@@ -92,10 +93,12 @@ template <typename T> class LinearAllocGuard {
|
||||
allocation_type_ = o.allocation_type_;
|
||||
ptr_ = o.ptr_;
|
||||
host_ptr_ = o.host_ptr_;
|
||||
size_ = o.size_;
|
||||
|
||||
o.allocation_type_ = LinearAllocs::noAlloc;
|
||||
o.ptr_ = nullptr;
|
||||
o.host_ptr_ = nullptr;
|
||||
o.size_ = 0;
|
||||
}
|
||||
|
||||
return *this;
|
||||
@@ -105,11 +108,13 @@ template <typename T> class LinearAllocGuard {
|
||||
|
||||
T* ptr() const { return ptr_; };
|
||||
T* host_ptr() const { return host_ptr_; }
|
||||
size_t size_bytes() const { return size_; }
|
||||
|
||||
private:
|
||||
LinearAllocs allocation_type_ = LinearAllocs::noAlloc;
|
||||
T* ptr_ = nullptr;
|
||||
T* host_ptr_ = nullptr;
|
||||
size_t size_ = 0;
|
||||
|
||||
void dealloc() {
|
||||
if (ptr_ == nullptr) {
|
||||
|
||||
@@ -0,0 +1,519 @@
|
||||
/*
|
||||
Copyright (c) 2025 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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
#define HIP_ENABLE_WARP_SYNC_BUILTINS
|
||||
#define HIP_ENABLE_EXTRA_WARP_SYNC_TYPES
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
#include <hip/hip_fp16.h>
|
||||
#include <limits>
|
||||
#include <cmath>
|
||||
#include <iostream>
|
||||
#include <ios>
|
||||
|
||||
#define MASK_SHIFT(x, n) \
|
||||
(x & (static_cast<uint64_t>(1) << n)) >> n
|
||||
|
||||
const unsigned long long Every5thBit = 0x1084210842108421;
|
||||
const unsigned long long Every9thBit = 0x8040201008040201;
|
||||
const unsigned long long Every5thBut9th = Every5thBit & ~Every9thBit;
|
||||
const unsigned long long AllThreads = ~0;
|
||||
static constexpr int kNumReduces = 5000;
|
||||
|
||||
inline __device__ bool deactivate_thread(const uint64_t* const active_masks) {
|
||||
const auto warp =
|
||||
cooperative_groups::tiled_partition(cooperative_groups::this_thread_block(), warpSize);
|
||||
const auto block = cooperative_groups::this_thread_block();
|
||||
const auto warps_per_block = (block.size() + warpSize - 1) / warpSize;
|
||||
const auto block_rank = (blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x;
|
||||
const auto idx = block_rank * warps_per_block + block.thread_rank() / warpSize;
|
||||
return !(active_masks[idx] & (static_cast<uint64_t>(1) << warp.thread_rank()));
|
||||
}
|
||||
|
||||
inline std::mt19937& GetRandomGenerator() {
|
||||
static std::mt19937 mt(std::random_device{}());
|
||||
return mt;
|
||||
}
|
||||
|
||||
template <typename T> inline T GenerateRandomInteger(const T min, const T max) {
|
||||
std::uniform_int_distribution<T> dist(min, max);
|
||||
return dist(GetRandomGenerator());
|
||||
}
|
||||
|
||||
template <typename T> inline T GenerateRandomReal(const T min, const T max) {
|
||||
std::uniform_real_distribution<T> dist(min, max);
|
||||
return dist(GetRandomGenerator());
|
||||
}
|
||||
|
||||
inline int generate_width(int warp_size) {
|
||||
int exponent = 0;
|
||||
while (warp_size >>= 1) {
|
||||
++exponent;
|
||||
}
|
||||
|
||||
return GENERATE_COPY(map([](int e) { return 1 << e; }, range(1, exponent + 1)));
|
||||
}
|
||||
|
||||
inline uint64_t get_active_mask(unsigned int warp_id, unsigned int warp_size) {
|
||||
uint64_t active_mask = 0;
|
||||
switch (warp_id % 5) {
|
||||
case 0: // even threads in the warp
|
||||
active_mask = 0xAAAAAAAAAAAAAAAA;
|
||||
break;
|
||||
case 1: // odd threads in the warp
|
||||
active_mask = 0x5555555555555555;
|
||||
break;
|
||||
case 2: // first half of the warp
|
||||
for (int i = 0; i < warp_size / 2; i++) {
|
||||
active_mask = active_mask | (static_cast<uint64_t>(1) << i);
|
||||
}
|
||||
break;
|
||||
case 3: // second half of the warp
|
||||
for (int i = warp_size / 2; i < warp_size; i++) {
|
||||
active_mask = active_mask | (static_cast<uint64_t>(1) << i);
|
||||
}
|
||||
break;
|
||||
case 4: // all threads
|
||||
active_mask = 0xFFFFFFFFFFFFFFFF;
|
||||
break;
|
||||
}
|
||||
return active_mask;
|
||||
}
|
||||
|
||||
template <typename T, std::enable_if_t<std::is_integral<T>::value, bool> = true>
|
||||
inline T expandPrecision(int X) { return X; }
|
||||
|
||||
template <typename T, std::enable_if_t<std::is_floating_point<T>::value, bool> = true>
|
||||
inline T expandPrecision(int X) {
|
||||
return X * 3.141592653589793115997963468544185161590576171875;
|
||||
}
|
||||
|
||||
template <typename T, std::enable_if_t<std::is_same<T, __half>::value, bool> = true>
|
||||
inline __half expandPrecision(int X) {
|
||||
return (__half)expandPrecision<float>(X);
|
||||
}
|
||||
|
||||
template <typename T, std::enable_if_t<std::is_same<T, __half2>::value, bool> = true>
|
||||
inline __half2 expandPrecision(int X) {
|
||||
__half H = expandPrecision<float>(X);
|
||||
return {H, H};
|
||||
}
|
||||
|
||||
template <typename T, std::enable_if_t<std::is_integral<T>::value, bool> = true>
|
||||
inline void expandPrecision(T* Array, int size) {
|
||||
(void)Array;
|
||||
(void)size;
|
||||
}
|
||||
|
||||
template <typename T, std::enable_if_t<std::is_floating_point<T>::value, bool> = true>
|
||||
inline void expandPrecision(T *Array, int size) {
|
||||
for (int i = 0; i != size; ++i) {
|
||||
Array[i] *= 3.141592653589793115997963468544185161590576171875;
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
inline void initializeInput(T *Input, int size) {
|
||||
int Values[] = {0, -1, 2, 3, 4, 5, -6, 7,
|
||||
8, -9, 10, 11, 12, 13, -14, 15,
|
||||
16, 17, -18, 19, 20, -21, 22, 23,
|
||||
24, 25, 26, -27, 28, 29, 30, 31,
|
||||
-32, 33, 34, 35, -36, 37, 38, -39,
|
||||
40, 41, 42, 43, -44, -45, 46, 47,
|
||||
48, 49, 50, -51, 52, 53, -54, 55,
|
||||
56, 57, -58, 59, 60, 61, 62, -63};
|
||||
|
||||
for (int i = 0; i != size; ++i) {
|
||||
Input[i] = expandPrecision<T>(Values[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
inline void initializeExpected(T *Expected, int *Values, int size) {
|
||||
for (int i = 0; i != size; ++i) {
|
||||
Expected[i] = expandPrecision<T>(Values[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
inline bool compareEqual(T X, T Y) { return X == Y; }
|
||||
|
||||
template <>
|
||||
inline bool compareEqual(__half X, __half Y) {
|
||||
return __half2float(X) == __half2float(Y);
|
||||
}
|
||||
|
||||
template <>
|
||||
inline bool compareEqual(__half2 X, __half2 Y) {
|
||||
return compareEqual(X.x, Y.x) && compareEqual(X.y, Y.y);
|
||||
}
|
||||
|
||||
inline bool compareMaskEqual(unsigned long long *Actual, unsigned long long *Expected,
|
||||
int i, int warpSize) {
|
||||
if (warpSize == 32)
|
||||
return (unsigned)Actual[i] == (unsigned)Expected[i];
|
||||
return Actual[i] == Expected[i];
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
inline T alignUp(T num, size_t n) {
|
||||
if (num % n == 0) {
|
||||
return num;
|
||||
}
|
||||
|
||||
return ((num + n - 1) / n) * n;
|
||||
}
|
||||
|
||||
template <class T>
|
||||
struct DistributionType {
|
||||
using type = std::uniform_int_distribution<T>;
|
||||
};
|
||||
|
||||
// there is no std::uniform_real_distribution for 'half' type, so we cast from
|
||||
// unsigned short, avoiding Nan and Infinity
|
||||
template <>
|
||||
struct DistributionType<__half> {
|
||||
using type = std::uniform_int_distribution<unsigned short>;
|
||||
};
|
||||
|
||||
template <>
|
||||
struct DistributionType<float> {
|
||||
using type = std::uniform_real_distribution<float>;
|
||||
};
|
||||
|
||||
template <>
|
||||
struct DistributionType<double> {
|
||||
using type = std::uniform_real_distribution<double>;
|
||||
};
|
||||
|
||||
|
||||
template <class T>
|
||||
struct MinOp {
|
||||
T operator()(const T& lhs, const T& rhs) const
|
||||
{
|
||||
return std::min(lhs, rhs);
|
||||
}
|
||||
};
|
||||
|
||||
template <class T>
|
||||
struct MaxOp {
|
||||
T operator()(const T& lhs, const T& rhs) const
|
||||
{
|
||||
return std::max(lhs, rhs);
|
||||
}
|
||||
};
|
||||
|
||||
template <class T>
|
||||
struct XorOp {
|
||||
__host__ __device__ T operator()(const T& lhs, const T& rhs)
|
||||
{
|
||||
return (!lhs) != (!rhs) == 1;
|
||||
}
|
||||
};
|
||||
|
||||
// typeid(T).name() does seem to return a very descriptive name for primitive types,
|
||||
// at least on clang, so we roll out an equivalent
|
||||
template<class T>
|
||||
const char* typeToString()
|
||||
{
|
||||
if (std::is_same<T, int>::value)
|
||||
return "int";
|
||||
if (std::is_same<T, unsigned int>::value)
|
||||
return "unsigned int";
|
||||
if (std::is_same<T, long long>::value)
|
||||
return "long long";
|
||||
if (std::is_same<T, unsigned long long>::value)
|
||||
return "unsigned long long";
|
||||
if (std::is_same<T, half>::value)
|
||||
return "half";
|
||||
if (std::is_same<T, float>::value)
|
||||
return "float";
|
||||
if (std::is_same<T, double>::value)
|
||||
return "double";
|
||||
|
||||
assert(false && "Missing conversion to string for type");
|
||||
return "";
|
||||
}
|
||||
|
||||
template<class T, template <typename> class Op>
|
||||
const char* opToString()
|
||||
{
|
||||
if constexpr (std::is_same<Op<T>, std::plus<T>>::value)
|
||||
return "add";
|
||||
else if constexpr (std::is_same<Op<T>, MinOp<T>>::value)
|
||||
return "min";
|
||||
else if constexpr (std::is_same<Op<T>, MaxOp<T>>::value)
|
||||
return "max";
|
||||
else if constexpr (std::is_same<Op<T>, std::logical_and<T>>::value)
|
||||
return "logical_and";
|
||||
else if constexpr (std::is_same<Op<T>, std::logical_or<T>>::value)
|
||||
return "logical_or";
|
||||
else if constexpr (std::is_same<Op<T>, XorOp<T>>::value)
|
||||
return "logical_xor";
|
||||
else {
|
||||
static_assert(std::is_void<T>::value, "Unsupported operator");
|
||||
return "";
|
||||
}
|
||||
}
|
||||
|
||||
template <class T, class Gen>
|
||||
void genRandomMasks(LinearAllocGuard<T>& d_buf,
|
||||
LinearAllocGuard<T>& buf,
|
||||
Gen& gen,
|
||||
int numItems)
|
||||
{
|
||||
// masks must be != 0, hence passing 1 as the 'a' distribution parameter
|
||||
std::uniform_int_distribution<unsigned long long> dist(1);
|
||||
int numBytes = numItems * sizeof(T);
|
||||
LinearAllocGuard<T> tmp(LinearAllocs::malloc, numBytes);
|
||||
LinearAllocGuard<T> d_tmp(LinearAllocs::hipMalloc, numBytes);
|
||||
|
||||
buf = std::move(tmp);
|
||||
d_buf = std::move(d_tmp);
|
||||
|
||||
for (int i = 0; i < numItems; i++) {
|
||||
T mask = dist(gen);
|
||||
|
||||
if (getWarpSize() == 32)
|
||||
mask &= 0xFFFFFFFF;
|
||||
|
||||
buf.ptr()[i] = mask;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_buf.ptr(), buf.ptr(), numBytes, hipMemcpyHostToDevice));
|
||||
}
|
||||
|
||||
// generates a random __half (instead of using uniform_real_distribution<float> casting to __half
|
||||
// which is problematic)
|
||||
// @expDist needs to be between [0-2^5-2]
|
||||
template <class Gen>
|
||||
__half genRandomHalf(std::uniform_int_distribution<unsigned short>& dist,
|
||||
Gen& gen)
|
||||
{
|
||||
__half_raw tmp;
|
||||
|
||||
tmp.x = dist(gen);
|
||||
// rewrite the exponent to force the number to be (-8<x<8) and at the same time avoid NaN or
|
||||
// infinity
|
||||
tmp.x &= 0xBBFF;
|
||||
return tmp;
|
||||
}
|
||||
|
||||
// generates a random buffer in buf, copies it to device memory in d_buf
|
||||
template <class T, class Dist, class Gen>
|
||||
void genRandomBuffers(LinearAllocGuard<T>& d_buf,
|
||||
LinearAllocGuard<T>& buf,
|
||||
Dist& dist,
|
||||
Gen& gen,
|
||||
int numItems)
|
||||
{
|
||||
int numBytes = numItems * sizeof(T);
|
||||
LinearAllocGuard<T> tmp(LinearAllocs::malloc, numBytes);
|
||||
LinearAllocGuard<T> d_tmp(LinearAllocs::hipMalloc, numBytes);
|
||||
|
||||
buf = std::move(tmp);
|
||||
d_buf = std::move(d_tmp);
|
||||
|
||||
for (int i = 0; i < numItems; i++)
|
||||
if constexpr (std::is_same<T, __half>::value)
|
||||
buf.ptr()[i] = genRandomHalf(dist, gen);
|
||||
else
|
||||
buf.ptr()[i] = dist(gen);
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_buf.ptr(), buf.ptr(), numBytes, hipMemcpyHostToDevice));
|
||||
}
|
||||
|
||||
// given an operation produces the expected result of the reduction
|
||||
// @mask indicates the lanes that will participate in the computation
|
||||
template <class T, class Op>
|
||||
T calculateExpected(const T* input, Op op, unsigned long long mask)
|
||||
{
|
||||
T result;
|
||||
int wavefrontSize = getWarpSize();
|
||||
|
||||
if (std::is_same<Op, std::plus<T>>::value) {
|
||||
T tmp[64] = { 0 };
|
||||
|
||||
for (int i = 0; i < wavefrontSize; i++) {
|
||||
if (mask & (1ul << i)) {
|
||||
tmp[i] = input[i];
|
||||
}
|
||||
}
|
||||
|
||||
for (int modulo = 2; modulo <= wavefrontSize; modulo *= 2) {
|
||||
for (int i = 0; i < wavefrontSize; i += modulo) {
|
||||
int j = i + modulo / 2;
|
||||
|
||||
if (j < wavefrontSize)
|
||||
tmp[i] += tmp[j];
|
||||
}
|
||||
}
|
||||
result = tmp[0];
|
||||
} else {
|
||||
bool initialized = false;
|
||||
|
||||
for (int i = 0; i < wavefrontSize; i++) {
|
||||
if (mask & (1ul << i)) {
|
||||
if (initialized)
|
||||
result = op(input[i], result);
|
||||
else {
|
||||
result = input[i];
|
||||
initialized = true;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
return result;
|
||||
}
|
||||
|
||||
template <class T>
|
||||
void printMismatch(const T& result, const T& expected, const T* input, unsigned long long mask)
|
||||
{
|
||||
std::ios init(NULL);
|
||||
|
||||
init.copyfmt(std::cout);
|
||||
std::cout << "\nMismatch\n";
|
||||
std::cout << "Mask: 0x" << std::hex << std::setfill('0') << std::setw(16) << mask << "\n";
|
||||
std::cout.copyfmt(init);
|
||||
|
||||
for (int i = 0; i < getWarpSize(); i++) {
|
||||
if ((1ul << i) & mask) {
|
||||
if constexpr (std::is_same<T, __half>::value)
|
||||
std::cout << "Lane " << i << ": " << __half2float(input[i]) << "\n";
|
||||
else
|
||||
std::cout << "Lane " << i << ": " << input[i] << "\n";
|
||||
}
|
||||
}
|
||||
|
||||
if constexpr (std::is_same<T, __half>::value) {
|
||||
std::cout << "Result: " << __half2float(result) << "\n";
|
||||
std::cout << "Expected: " << __half2float(expected) << "\n";
|
||||
} else {
|
||||
std::cout << "Result: " << result << "\n";
|
||||
std::cout << "Expected: " << expected << "\n";
|
||||
}
|
||||
}
|
||||
|
||||
template <class T>
|
||||
void compareFloatingPoint(const T& result, const T& expected, unsigned long long mask, const T* input)
|
||||
{
|
||||
using namespace Catch::Matchers;
|
||||
if constexpr (std::is_same<T, __half>::value) {
|
||||
float resultFloat = __half2float(result);
|
||||
float expectedFloat = __half2float(expected);
|
||||
float absDifference = fabs(resultFloat - expectedFloat);
|
||||
float relativeEpsilon = 0.1 * fmax(resultFloat, expectedFloat);
|
||||
float eps = 0.01f;
|
||||
|
||||
REQUIRE(!__hisnan(result));
|
||||
REQUIRE(!__hisinf(result));
|
||||
|
||||
if (relativeEpsilon > eps) {
|
||||
if (absDifference > 0.0001) {
|
||||
if (absDifference >= eps * fabs(fmax(resultFloat, expectedFloat))) {
|
||||
printMismatch(result, expected, input, mask);
|
||||
std::cout << "Relative epsilon: " << relativeEpsilon << "\n";
|
||||
std::cout << "Difference: " << absDifference << "\n";
|
||||
}
|
||||
}
|
||||
|
||||
REQUIRE_THAT(__half2float(resultFloat), WithinRel(expectedFloat, eps));
|
||||
}
|
||||
} else {
|
||||
// for float or double, also lossy in terms of precision
|
||||
T absDifference = fabs(result - expected);
|
||||
T relativeEpsilon = 0.1 * fmax(result, expected);
|
||||
T eps = 0.01;
|
||||
|
||||
if (relativeEpsilon > eps) {
|
||||
if (absDifference > 0.0001) {
|
||||
if (absDifference >= eps * fabs(fmax(result, expected))) {
|
||||
printMismatch(result, expected, input, mask);
|
||||
std::cout << "Relative epsilon: " << relativeEpsilon << "\n";
|
||||
std::cout << "Difference: " << absDifference << "\n";
|
||||
}
|
||||
|
||||
REQUIRE_THAT(result, WithinRel(expected, eps));
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// @tparam Reduce a functor; abstracts away kernel dispatching
|
||||
// (via hiprtc or normal execution)
|
||||
template <class T, class Reduce, template <typename> class Op>
|
||||
void runTestReduce(int iteration, Reduce reduce)
|
||||
{
|
||||
using namespace Catch::Matchers;
|
||||
using distribution = typename DistributionType<T>::type;
|
||||
unsigned int wavefrontSize = getWarpSize();
|
||||
// one result per reduce per thread to be checked
|
||||
LinearAllocGuard<T> d_output(LinearAllocs::hipMalloc, kNumReduces * wavefrontSize * sizeof(T));
|
||||
LinearAllocGuard<T> output(LinearAllocs::malloc, kNumReduces * wavefrontSize * sizeof(T));
|
||||
std::mt19937_64 gen(iteration);
|
||||
// for float16, we generate any random unsigned short, but cap the exponent later on
|
||||
// to keep it in the range (-8.0..8.0) (just to avoid overflows)
|
||||
// On the rest of the types, just use a bigger reduced range of numbers to avoid overflows too
|
||||
T a = std::is_same<T, half>::value? std::numeric_limits<unsigned short>::lowest() : -1023;
|
||||
T b = std::is_same<T, half>::value? std::numeric_limits<unsigned short>::max() : 1023;
|
||||
distribution dist(a, b);
|
||||
LinearAllocGuard<T> input, d_input;
|
||||
LinearAllocGuard<unsigned long long> masks, d_masks;
|
||||
Op<T> op;
|
||||
int numReduce = 0;
|
||||
|
||||
genRandomBuffers(d_input, input, dist, gen, kNumReduces * wavefrontSize);
|
||||
genRandomMasks(d_masks, masks, gen, kNumReduces);
|
||||
reduce(d_output.ptr(), d_input.ptr(), d_masks.ptr(), kNumReduces, op);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
HIP_CHECK(hipMemcpy(output.ptr(), d_output.ptr(), d_output.size_bytes(), hipMemcpyDeviceToHost));
|
||||
|
||||
while (numReduce < kNumReduces) {
|
||||
T expected = calculateExpected<T>(input.ptr(), op, masks.ptr()[numReduce]);
|
||||
int lane = 0;
|
||||
|
||||
while (lane < wavefrontSize) {
|
||||
auto result = output.ptr()[numReduce * wavefrontSize + lane];
|
||||
unsigned long long mask = masks.ptr()[numReduce];
|
||||
|
||||
if ((1ul << lane) & mask) {
|
||||
if constexpr (std::is_integral<T>::value || std::is_same<Op<T>, MinOp<T>>::value ||
|
||||
std::is_same<Op<T>, MaxOp<T>>::value) {
|
||||
// for integral types or min/max the result should match exactly
|
||||
if constexpr (std::is_same<T, __half>::value)
|
||||
REQUIRE(__half2float(result) == __half2float(expected));
|
||||
else {
|
||||
if (result != expected) {
|
||||
printMismatch(result, expected, input.ptr(), mask);
|
||||
REQUIRE(result == expected);
|
||||
}
|
||||
}
|
||||
} else
|
||||
compareFloatingPoint(result, expected, mask, input.ptr());
|
||||
|
||||
}
|
||||
lane++;
|
||||
}
|
||||
numReduce++;
|
||||
}
|
||||
}
|
||||
@@ -23,4 +23,5 @@ add_subdirectory(memcpy)
|
||||
add_subdirectory(kernelLaunch)
|
||||
add_subdirectory(stream)
|
||||
add_subdirectory(event)
|
||||
add_subdirectory(warpSync)
|
||||
add_subdirectory(example)
|
||||
|
||||
@@ -0,0 +1,31 @@
|
||||
# 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.
|
||||
|
||||
# As CUDA use 32-bit sync masks and not 64-bit this has been disabled on Nvidia
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set(TEST_SRC
|
||||
warpSync.cc
|
||||
)
|
||||
endif()
|
||||
|
||||
hip_add_exe_to_target(NAME WarpSyncPerformance
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests
|
||||
COMPILE_OPTIONS -std=c++17)
|
||||
@@ -0,0 +1,395 @@
|
||||
/*
|
||||
Copyright (c) 2025 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.
|
||||
*/
|
||||
#define HIP_ENABLE_WARP_SYNC_BUILTINS
|
||||
#define HIP_ENABLE_EXTRA_WARP_SYNC_TYPES
|
||||
|
||||
#include "warp_common.hh"
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_fp16.h>
|
||||
#include <hip_test_common.hh>
|
||||
#include <performance_common.hh>
|
||||
#include <hip/amd_detail/amd_hip_atomic.h>
|
||||
#include <resource_guards.hh>
|
||||
#include <cmd_options.hh>
|
||||
#include <algorithm>
|
||||
#include <type_traits>
|
||||
#include <limits>
|
||||
#include <map>
|
||||
|
||||
/**
|
||||
* @addtogroup __reduce_op_sync __reduce_op_sync
|
||||
* @{
|
||||
* @ingroup WarpSyncPerformance
|
||||
* __reduce_op_sync(MaskT mask, T val)
|
||||
* Reduces the val as per the lanes described in mask and calculates the
|
||||
* aggregated result
|
||||
*/
|
||||
|
||||
static constexpr int kBlockDim = 1024;
|
||||
|
||||
template <class T>
|
||||
struct AtomicAddOp {
|
||||
__device__ T operator()(T* lhs, const T& rhs)
|
||||
{
|
||||
return atomicAdd(lhs, rhs);
|
||||
}
|
||||
};
|
||||
|
||||
template <class T>
|
||||
struct AtomicMinOp {
|
||||
__device__ T operator()(T* lhs, const T& rhs)
|
||||
{
|
||||
return atomicMin(lhs, rhs);
|
||||
}
|
||||
};
|
||||
|
||||
template <class T>
|
||||
struct AtomicMaxOp {
|
||||
__device__ T operator()(T* lhs, const T& rhs)
|
||||
{
|
||||
return atomicMax(lhs, rhs);
|
||||
}
|
||||
};
|
||||
|
||||
template <class T>
|
||||
struct AtomicAndOp {
|
||||
__device__ T operator()(T* lhs, const T& rhs)
|
||||
{
|
||||
return atomicAnd(lhs, rhs);
|
||||
}
|
||||
};
|
||||
|
||||
template <class T>
|
||||
struct AtomicOrOp {
|
||||
__device__ T operator()(T* lhs, const T& rhs)
|
||||
{
|
||||
return atomicOr(lhs, rhs);
|
||||
}
|
||||
};
|
||||
|
||||
template <class T>
|
||||
struct AtomicXorOp {
|
||||
__device__ T operator()(T* lhs, const T& rhs)
|
||||
{
|
||||
return atomicXor(lhs, rhs);
|
||||
}
|
||||
};
|
||||
|
||||
// uses atomics to reduce the whole warp; depending on the mask our reduce should be faster
|
||||
// @output to store the result, one per warp
|
||||
// @numItems must be a multiple of warpSize
|
||||
template <class T, template <typename> class Op>
|
||||
__global__ void reduceAllAtomics(T* __restrict__ output, const T* __restrict__ input, unsigned long long mask)
|
||||
{
|
||||
int idx = threadIdx.x + blockIdx.x * kBlockDim;
|
||||
__shared__ T result[kBlockDim / warpSize]; // one per warp
|
||||
Op<T> op;
|
||||
int numWarp = threadIdx.x / warpSize;
|
||||
|
||||
// initialize result[numWarp] to the "identity" element for Op
|
||||
if constexpr (std::is_same<Op<T>, AtomicMinOp<T>>::value)
|
||||
result[numWarp] = std::numeric_limits<T>::max();
|
||||
else if constexpr (std::is_same<Op<T>, AtomicMaxOp<T>>::value)
|
||||
result[numWarp] = std::numeric_limits<T>::lowest();
|
||||
else if constexpr (std::is_same<Op<T>, AtomicAndOp<T>>::value)
|
||||
result[numWarp] = 1;
|
||||
else
|
||||
result[numWarp] = 0;
|
||||
|
||||
__syncthreads();
|
||||
|
||||
if (mask & (1ul << __ockl_lane_u32()))
|
||||
op(&result[numWarp], input[idx]);
|
||||
|
||||
__syncthreads();
|
||||
|
||||
if (__ockl_lane_u32() == 0)
|
||||
output[idx / warpSize] = result[numWarp];
|
||||
}
|
||||
|
||||
template <class T, template<typename> class Op>
|
||||
__global__ void reduceOpSync(T* __restrict__ output, const T* __restrict__ input, unsigned long long mask)
|
||||
{
|
||||
int idx = threadIdx.x + blockIdx.x * kBlockDim;
|
||||
T result;
|
||||
|
||||
if (mask & (1ul << __ockl_lane_u32())) {
|
||||
if constexpr (std::is_same<Op<T>, std::plus<T>>::value)
|
||||
result = __reduce_add_sync(mask, input[idx]);
|
||||
else if constexpr (std::is_same<Op<T>, MinOp<T>>::value)
|
||||
result = __reduce_min_sync(mask, input[idx]);
|
||||
else if constexpr (std::is_same<Op<T>, MaxOp<T>>::value)
|
||||
result = __reduce_max_sync(mask, input[idx]);
|
||||
else if constexpr (std::is_same<Op<T>, std::logical_and<T>>::value)
|
||||
result = __reduce_and_sync(mask, input[idx]);
|
||||
else if constexpr (std::is_same<Op<T>, std::logical_or<T>>::value)
|
||||
result = __reduce_or_sync(mask, input[idx]);
|
||||
else if constexpr (std::is_same<Op<T>, XorOp<T>>::value)
|
||||
result = __reduce_xor_sync(mask, input[idx]);
|
||||
else
|
||||
static_assert(std::is_void<T>::value, "Unsupported operator");
|
||||
|
||||
if (__ockl_activelane_u32() == 0)
|
||||
output[idx / warpSize] = result;
|
||||
}
|
||||
}
|
||||
|
||||
template <class T, template <typename> class Op>
|
||||
class AtomicBenchmark : public Benchmark<AtomicBenchmark<T, Op>> {
|
||||
public:
|
||||
void operator()(T* output, const T* input, int numItems, unsigned long long mask)
|
||||
{
|
||||
dim3 blockDim = { kBlockDim };
|
||||
dim3 gridDim = { static_cast<uint32_t>(std::ceil(numItems / static_cast<float>(blockDim.x))) };
|
||||
|
||||
TIMED_SECTION(kTimerTypeEvent) {
|
||||
if constexpr (std::is_same<Op<T>, std::plus<T>>::value)
|
||||
reduceAllAtomics<T, AtomicAddOp><<<gridDim, blockDim>>>(output, input, mask);
|
||||
else if constexpr (std::is_same<Op<T>, MinOp<T>>::value)
|
||||
reduceAllAtomics<T, AtomicMinOp><<<gridDim, blockDim>>>(output, input, mask);
|
||||
else if constexpr (std::is_same<Op<T>, MaxOp<T>>::value)
|
||||
reduceAllAtomics<T, AtomicMaxOp><<<gridDim, blockDim>>>(output, input, mask);
|
||||
else if constexpr (std::is_same<Op<T>, std::logical_and<T>>::value)
|
||||
reduceAllAtomics<T, AtomicAndOp><<<gridDim, blockDim>>>(output, input, mask);
|
||||
else if constexpr (std::is_same<Op<T>, std::logical_or<T>>::value)
|
||||
reduceAllAtomics<T, AtomicOrOp><<<gridDim, blockDim>>>(output, input, mask);
|
||||
else if constexpr (std::is_same<Op<T>, XorOp<T>>::value)
|
||||
reduceAllAtomics<T, AtomicXorOp><<<gridDim, blockDim>>>(output, input, mask);
|
||||
else
|
||||
static_assert(std::is_void<T>::value, "Unsupported operator");
|
||||
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
template <class T, template <typename> class Op>
|
||||
class ReduceSyncBenchmark : public Benchmark<ReduceSyncBenchmark<T, Op>> {
|
||||
public:
|
||||
void operator()(T* output, T* input, int numItems, unsigned long long mask)
|
||||
{
|
||||
dim3 blockDim = { kBlockDim };
|
||||
dim3 gridDim = { static_cast<uint32_t>(std::ceil(numItems / static_cast<float>(blockDim.x))) };
|
||||
|
||||
|
||||
TIMED_SECTION(kTimerTypeEvent) {
|
||||
reduceOpSync<T, Op><<<gridDim, blockDim>>>(output, input, mask);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
template <class T, template <typename> class Op>
|
||||
void checkResults(T* d_atomicsResult, T* d_reduceResult, size_t numBytes, unsigned long long mask)
|
||||
{
|
||||
using namespace Catch::Matchers;
|
||||
LinearAllocGuard<T> outputAtomic(LinearAllocs::malloc, numBytes);
|
||||
LinearAllocGuard<T> outputReduce(LinearAllocs::malloc, numBytes);
|
||||
bool memcmpResult = std::memcmp(outputAtomic.ptr(), outputReduce.ptr(), numBytes);
|
||||
|
||||
assert(numBytes % sizeof(T) == 0 && "numBytes needs to be a multiple of sizeof(T)");
|
||||
HIP_CHECK(hipMemcpy(outputAtomic.ptr(), d_atomicsResult, numBytes, hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipMemcpy(outputReduce.ptr(), d_reduceResult, numBytes, hipMemcpyDeviceToHost));
|
||||
|
||||
if (memcmpResult) {
|
||||
for (int i = 0; i < numBytes / sizeof(T); i++) {
|
||||
auto& atomicResult = outputAtomic.ptr()[i];
|
||||
auto& reduceResult = outputReduce.ptr()[i];
|
||||
|
||||
if constexpr (std::is_integral<T>::value || std::is_same<Op<T>, MinOp<T>>::value ||
|
||||
std::is_same<Op<T>, MaxOp<T>>::value)
|
||||
// for integral types or min/max the result should match exactly
|
||||
REQUIRE(atomicResult == reduceResult);
|
||||
else
|
||||
// floating point types or operations which are lossy in terms of precision
|
||||
REQUIRE_THAT(reduceResult, WithinRel(atomicResult));
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
template <class T, template <typename> class Op>
|
||||
struct IsLogicalOp {
|
||||
static constexpr bool value = false;
|
||||
};
|
||||
|
||||
template <class T>
|
||||
struct IsLogicalOp<T, std::logical_and> {
|
||||
static constexpr bool value = true;
|
||||
};
|
||||
|
||||
template <class T>
|
||||
struct IsLogicalOp<T, std::logical_or> {
|
||||
static constexpr bool value = true;
|
||||
};
|
||||
|
||||
template <class T>
|
||||
struct IsLogicalOp<T, XorOp> {
|
||||
static constexpr bool value = true;
|
||||
};
|
||||
|
||||
// Neither long long or fp16 have atomic operations. In those cases
|
||||
// we only benchmark reduce sync operations, we cannot compare with native atomics
|
||||
template <class T>
|
||||
struct HasAtomicOps {
|
||||
static constexpr bool value = true;
|
||||
};
|
||||
|
||||
template <>
|
||||
struct HasAtomicOps<half> {
|
||||
static constexpr bool value = false;
|
||||
};
|
||||
|
||||
template <>
|
||||
struct HasAtomicOps<long long> {
|
||||
static constexpr bool value = false;
|
||||
};
|
||||
|
||||
template <class T, template <typename> class Op>
|
||||
struct ReduceBenchmark {
|
||||
void Run()
|
||||
{
|
||||
static constexpr int numMasks = 6;
|
||||
using distribution = typename DistributionType<T>::type;
|
||||
ReduceSyncBenchmark<T, Op> benchmarkReduce;
|
||||
uint64_t inputSize = cmd_options.reduce_input_size * 1_MB;
|
||||
int numItems = inputSize / sizeof(T);
|
||||
int wavefrontSize = getWarpSize();
|
||||
int outputNumBytes = inputSize / wavefrontSize;
|
||||
LinearAllocGuard<T> input(LinearAllocs::malloc, inputSize);
|
||||
LinearAllocGuard<T> d_input(LinearAllocs::hipMalloc, inputSize);
|
||||
LinearAllocGuard<T> d_outputsAtomic[numMasks];
|
||||
LinearAllocGuard<T> d_outputsReduce[numMasks];
|
||||
LinearAllocGuard<T>* d_outputAtomic = &d_outputsAtomic[0];
|
||||
LinearAllocGuard<T>* d_outputReduce = &d_outputsReduce[0];
|
||||
std::mt19937_64 gen(123);
|
||||
distribution dist;
|
||||
int halfWaveSize = wavefrontSize / 2;
|
||||
unsigned long long halfBitsOn = (1ul << (wavefrontSize / 2)) - 1;
|
||||
unsigned long long fullMask = -1ul,
|
||||
halfHighBitsOn = halfBitsOn << halfWaveSize,
|
||||
high16BitsOn = halfBitsOn << (wavefrontSize - 16),
|
||||
high8BitsOn = halfBitsOn << (wavefrontSize - 8),
|
||||
high4BitsOn = halfBitsOn << (wavefrontSize - 4),
|
||||
allButOne = -1 & ~1;
|
||||
const char* typeStr = typeToString<T>();
|
||||
const char* opStr = opToString<T, Op>();
|
||||
std::map<std::string, unsigned long long> masks;
|
||||
std::pair<std::string, unsigned long long> masksPairs[] = { { "full mask", fullMask },
|
||||
{ "high order 32 bits on", halfHighBitsOn },
|
||||
{ "high order 16 bits on", high16BitsOn },
|
||||
{ "high order 8 bits on", high8BitsOn },
|
||||
{ "high order 4 bits on", high4BitsOn },
|
||||
{ "all but one", allButOne } };
|
||||
int pos = 0, numMask = 0;
|
||||
|
||||
for (auto& mask : masksPairs) {
|
||||
// don't use 'halfHighBitsOn' on warp size 32; it's the same as high16BitsOn
|
||||
if (wavefrontSize != 32 || mask.second != halfHighBitsOn) {
|
||||
masks.emplace(std::to_string(numMask) + " - " + mask.first, wavefrontSize == 64? mask.second : mask.second & 0xFFFFFFFF);
|
||||
numMask++;
|
||||
}
|
||||
}
|
||||
|
||||
// avoid generating values different than 1 or 0 for logical operators;
|
||||
// otherwise the atomic version of the kernels would produce different results as
|
||||
// atomicAnd/Or() are bitwise operations, not logical
|
||||
if constexpr (IsLogicalOp<T, Op>::value)
|
||||
dist = distribution(0, 1);
|
||||
|
||||
for (int i = 0; i < numItems; i++) {
|
||||
input.ptr()[i] = dist(gen);
|
||||
}
|
||||
|
||||
for (auto& buffer : d_outputsAtomic) {
|
||||
buffer = LinearAllocGuard<T>(LinearAllocs::hipMalloc, outputNumBytes);
|
||||
}
|
||||
|
||||
for (auto& buffer : d_outputsReduce) {
|
||||
buffer = LinearAllocGuard<T>(LinearAllocs::hipMalloc, outputNumBytes);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_input.ptr(), input.ptr(), inputSize, hipMemcpyHostToDevice));
|
||||
|
||||
if constexpr (HasAtomicOps<T>::value) {
|
||||
AtomicBenchmark<T, Op> benchmarkAtomics;
|
||||
|
||||
printf("\n--- atomics %s %s---\n", opStr, typeStr);
|
||||
|
||||
for (auto& mask : masks) {
|
||||
printf("%s %llx\n", mask.first.c_str(), mask.second);
|
||||
benchmarkAtomics.Run((d_outputAtomic++)->ptr(), d_input.ptr(), numItems, mask.second);
|
||||
}
|
||||
}
|
||||
|
||||
printf("\n--- reduce %s %s--- \n", opStr, typeStr);
|
||||
|
||||
for (const auto& mask : masks) {
|
||||
printf("%s %llx\n", mask.first.c_str(), mask.second);
|
||||
benchmarkReduce.Run((d_outputReduce++)->ptr(), d_input.ptr(), numItems, mask.second);
|
||||
}
|
||||
|
||||
printf("\n");
|
||||
|
||||
if constexpr (HasAtomicOps<T>::value) {
|
||||
printf("Checking results...\n");
|
||||
|
||||
for (const auto& mask : masks) {
|
||||
checkResults<T, Op>(d_outputsAtomic[pos].ptr(), d_outputsReduce[pos].ptr(), outputNumBytes, mask.second);
|
||||
pos++;
|
||||
}
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
TEMPLATE_TEST_CASE("Performance_Reduce_Sync_Add", "", int, unsigned int, unsigned long long,
|
||||
long long, float, half, double) {
|
||||
ReduceBenchmark<TestType, std::plus> benchmark;
|
||||
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Performance_Reduce_Sync_Min", "", int, unsigned int, unsigned long long, long long, float, half, double) {
|
||||
ReduceBenchmark<TestType, MinOp> benchmark;
|
||||
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Performance_Reduce_Sync_Max", "", int, unsigned int, unsigned long long, long long, float, half, double) {
|
||||
ReduceBenchmark<TestType, MaxOp> benchmark;
|
||||
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Performance_Reduce_Sync_And", "", int, unsigned int, unsigned long long, long long) {
|
||||
ReduceBenchmark<TestType, std::logical_and> benchmark;
|
||||
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Performance_Reduce_Sync_Or", "", int, unsigned int, unsigned long long, long long) {
|
||||
ReduceBenchmark<TestType, std::logical_or> benchmark;
|
||||
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Performance_Reduce_Sync_Xor", "", int, unsigned int, unsigned long long, long long) {
|
||||
ReduceBenchmark<TestType, XorOp> benchmark;
|
||||
|
||||
benchmark.Run();
|
||||
}
|
||||
@@ -16,6 +16,7 @@ set(AMD_TEST_SRC
|
||||
linker.cc
|
||||
shfl.cc
|
||||
shfl_sync.cc
|
||||
rtc_reduce.cc
|
||||
stdheaders.cc
|
||||
hiprtc_MathConstants_HeaderTst.cc
|
||||
hiprtc_VectorTypes_HeaderTst.cc
|
||||
|
||||
@@ -0,0 +1,215 @@
|
||||
/*
|
||||
Copyright (c) 2025 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.
|
||||
*/
|
||||
#define HIP_ENABLE_WARP_SYNC_BUILTINS
|
||||
#define HIP_ENABLE_EXTRA_WARP_SYNC_TYPES
|
||||
|
||||
#include "warp_common.hh"
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <tuple>
|
||||
#include <cmd_options.hh>
|
||||
#include <functional>
|
||||
#include <algorithm>
|
||||
|
||||
#define NELEMS(array) (sizeof(array) / sizeof(array[0]))
|
||||
|
||||
// compiles the program, reusing the same compiling session for all the types
|
||||
// (as opposed as calling the rtc compiler for each of the types)
|
||||
template <template <typename> class Op, class T, typename... Types>
|
||||
void compileProgram(hiprtcProgram& prog, const std::tuple<T, Types...>&)
|
||||
{
|
||||
std::string scalarName, intrinsicName, expression;
|
||||
std::tuple<Types...> remainingTypes;
|
||||
|
||||
expression = std::string("reduceRtcKernel<") + typeToString<T>() + ", unsigned long long>";
|
||||
HIPRTC_CHECK(hiprtcAddNameExpression(prog, expression.c_str()));
|
||||
compileProgram<Op>(prog, remainingTypes);
|
||||
}
|
||||
|
||||
template <class T, class MaskType, template <typename> class Op>
|
||||
void runRtcReduceOp(hiprtcProgram& prog, T* output, const T* input, const MaskType* masks, int numReduces, Op<T>)
|
||||
{
|
||||
unsigned int wavefrontSize = getWarpSize();
|
||||
const char* loweredName;
|
||||
hipFunction_t kernel;
|
||||
hipModule_t module;
|
||||
struct {
|
||||
const T* d_output;
|
||||
const T* d_input;
|
||||
const MaskType* d_masks;
|
||||
int numReduces;
|
||||
} args {output, input, masks, numReduces};
|
||||
int size = 4;
|
||||
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END};
|
||||
std::vector<char> code;
|
||||
size_t codeSize;
|
||||
std::string expression = std::string("reduceRtcKernel<") + typeToString<T>() + ", unsigned long long>";
|
||||
dim3 grdDim { 1u };
|
||||
dim3 blkDim { wavefrontSize };
|
||||
|
||||
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
|
||||
code.resize(codeSize);
|
||||
HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
|
||||
HIP_CHECK(hipModuleLoadData(&module, code.data()));
|
||||
HIPRTC_CHECK(hiprtcGetLoweredName(prog, expression.c_str(), &loweredName));
|
||||
HIP_CHECK(hipModuleGetFunction(&kernel, module, loweredName));
|
||||
HIP_CHECK(hipModuleLaunchKernel(kernel, grdDim.x, grdDim.y, grdDim.z, blkDim.x, blkDim.y, blkDim.z, 0, 0, nullptr, config));
|
||||
HIP_CHECK(hipModuleUnload(module));
|
||||
}
|
||||
|
||||
template <template <typename> class Op, class Type = void>
|
||||
void runTestReduceForTypes(hiprtcProgram&, const std::tuple<>)
|
||||
{
|
||||
}
|
||||
|
||||
template <template <typename> class Op, class T, typename... Types>
|
||||
void runTestReduceForTypes(hiprtcProgram& prog, const std::tuple<T, Types...>)
|
||||
{
|
||||
std::tuple<Types...> remainingTypes;
|
||||
int iteration = 0;
|
||||
|
||||
auto reduceFunc = [&prog](T* d_output, const T* d_input, const unsigned long long* d_masks, int numReduces, Op<T> op) {
|
||||
runRtcReduceOp(prog, d_output, d_input, d_masks, numReduces, op);
|
||||
};
|
||||
|
||||
while (iteration < cmd_options.reduce_iterations) {
|
||||
runTestReduce<T, decltype(reduceFunc), Op>(iteration, reduceFunc);
|
||||
iteration++;
|
||||
|
||||
if (cmd_options.reduce_iterations != 1) {
|
||||
std::cout << "\rIteration: " << iteration;
|
||||
std::flush(std::cout);
|
||||
}
|
||||
}
|
||||
|
||||
runTestReduceForTypes<Op>(prog, remainingTypes);
|
||||
}
|
||||
|
||||
template<class T, template <typename> class Op>
|
||||
void opToString(std::string& scalarName, std::string& intrinsicName)
|
||||
{
|
||||
if constexpr (std::is_same<Op<T>, std::plus<T>>::value) {
|
||||
scalarName = "std::plus";
|
||||
intrinsicName = "__reduce_add_sync";
|
||||
} else if constexpr (std::is_same<Op<T>, MinOp<T>>::value) {
|
||||
scalarName = "MinOp";
|
||||
intrinsicName = "__reduce_min_sync";
|
||||
} else if constexpr (std::is_same<Op<T>, MaxOp<T>>::value) {
|
||||
scalarName = "MaxOp";
|
||||
intrinsicName = "__reduce_max_sync";
|
||||
} else if constexpr (std::is_same<Op<T>, std::logical_and<T>>::value) {
|
||||
scalarName = "std::logical_and";
|
||||
intrinsicName = "__reduce_and_sync";
|
||||
} else if constexpr (std::is_same<Op<T>, std::logical_or<T>>::value) {
|
||||
scalarName = "std::logical_or";
|
||||
intrinsicName = "__reduce_or_sync";
|
||||
} else if constexpr (std::is_same<Op<T>, XorOp<T>>::value) {
|
||||
scalarName = "LogicalXor";
|
||||
intrinsicName = "__reduce_xor_sync";
|
||||
} else
|
||||
static_assert(std::is_void<T>::value, "Unexpected operator");
|
||||
}
|
||||
|
||||
template <template <typename> class Op, class T = void>
|
||||
void compileProgram(hiprtcProgram& prog, const std::tuple<>&)
|
||||
{
|
||||
size_t logSize;
|
||||
std::string scalarName, intrinsicName;
|
||||
hiprtcResult compileResult;
|
||||
const char* options[] = { "-DHIP_ENABLE_WARP_SYNC_BUILTINS", "-DHIP_ENABLE_EXTRA_WARP_SYNC_TYPES" };
|
||||
|
||||
opToString<int, Op>(scalarName, intrinsicName);
|
||||
compileResult = hiprtcResult {hiprtcCompileProgram(prog, NELEMS(options), options)};
|
||||
HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
|
||||
|
||||
if (compileResult != HIPRTC_SUCCESS || logSize > 0) {
|
||||
std::string log(logSize, '\0');
|
||||
|
||||
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
|
||||
std::cerr << "Runtime compilation failed or contained warnings for operator: "
|
||||
<< scalarName
|
||||
<< " associated reduce function: "
|
||||
<< intrinsicName
|
||||
<< "\n";
|
||||
std::cerr << log << '\n';
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
template <template <typename> class Op, typename... Types>
|
||||
void runAndCompileTest(const std::tuple<Types...> types)
|
||||
{
|
||||
std::string scalarName, intrinsicName, kernelStr;
|
||||
hiprtcProgram prog;
|
||||
|
||||
opToString<int, Op>(scalarName, intrinsicName);
|
||||
kernelStr = R"(
|
||||
template <class T, class MaskType>
|
||||
__global__ void reduceRtcKernel(T* output, const T* input, const MaskType* masks, int numReduces)
|
||||
{
|
||||
int tid = threadIdx.x;
|
||||
|
||||
for (int i = 0; i < numReduces; i++) {
|
||||
if (masks[i] & (1ul << tid)) {
|
||||
// call the operator only if the lane is mentioned in the mask
|
||||
T& result = output[warpSize * i + tid];
|
||||
result = )" + intrinsicName + R"((masks[i], input[tid]);
|
||||
}
|
||||
}
|
||||
})";
|
||||
|
||||
HIPRTC_CHECK(hiprtcCreateProgram(&prog,
|
||||
kernelStr.c_str(),
|
||||
"warp_reduce.hip",
|
||||
0,
|
||||
nullptr,
|
||||
nullptr));
|
||||
compileProgram<Op>(prog, types);
|
||||
runTestReduceForTypes<Op>(prog, types);
|
||||
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_Rtc_ReduceRandom") {
|
||||
const std::tuple<int, unsigned int, long long, unsigned long long, float, half, double> allTypes;
|
||||
const std::tuple<int, unsigned int, long long, unsigned long long> integralTypes;
|
||||
|
||||
SECTION("add") {
|
||||
runAndCompileTest<std::plus>(allTypes);
|
||||
}
|
||||
|
||||
SECTION("min") {
|
||||
runAndCompileTest<MinOp>(allTypes);
|
||||
}
|
||||
|
||||
SECTION("max") {
|
||||
runAndCompileTest<MaxOp>(allTypes);
|
||||
}
|
||||
|
||||
SECTION("and") {
|
||||
runAndCompileTest<std::logical_and>(integralTypes);
|
||||
}
|
||||
|
||||
SECTION("or") {
|
||||
runAndCompileTest<std::logical_or>(integralTypes);
|
||||
}
|
||||
|
||||
SECTION("xor") {
|
||||
runAndCompileTest<XorOp>(integralTypes);
|
||||
}
|
||||
}
|
||||
@@ -18,6 +18,7 @@ if(HIP_PLATFORM MATCHES "amd")
|
||||
warp_shfl_xor.cc
|
||||
warp_shfl_up.cc
|
||||
warp_shfl_down.cc
|
||||
warp_reduce.cc
|
||||
hipShflUpDownTest.cc
|
||||
hipShflTests.cc
|
||||
)
|
||||
|
||||
@@ -1,178 +0,0 @@
|
||||
/*
|
||||
Copyright (c) 2023 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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
#define HIP_ENABLE_WARP_SYNC_BUILTINS
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
#include <hip/hip_fp16.h>
|
||||
|
||||
#define MASK_SHIFT(x, n) \
|
||||
(x & (static_cast<uint64_t>(1) << n)) >> n
|
||||
|
||||
const unsigned long long Every5thBit = 0x1084210842108421;
|
||||
const unsigned long long Every9thBit = 0x8040201008040201;
|
||||
const unsigned long long Every5thBut9th = Every5thBit & ~Every9thBit;
|
||||
const unsigned long long AllThreads = ~0;
|
||||
|
||||
inline __device__ bool deactivate_thread(const uint64_t* const active_masks) {
|
||||
const auto warp =
|
||||
cooperative_groups::tiled_partition(cooperative_groups::this_thread_block(), warpSize);
|
||||
const auto block = cooperative_groups::this_thread_block();
|
||||
const auto warps_per_block = (block.size() + warpSize - 1) / warpSize;
|
||||
const auto block_rank = (blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x;
|
||||
const auto idx = block_rank * warps_per_block + block.thread_rank() / warpSize;
|
||||
return !(active_masks[idx] & (static_cast<uint64_t>(1) << warp.thread_rank()));
|
||||
}
|
||||
|
||||
inline std::mt19937& GetRandomGenerator() {
|
||||
static std::mt19937 mt(std::random_device{}());
|
||||
return mt;
|
||||
}
|
||||
|
||||
template <typename T> inline T GenerateRandomInteger(const T min, const T max) {
|
||||
std::uniform_int_distribution<T> dist(min, max);
|
||||
return dist(GetRandomGenerator());
|
||||
}
|
||||
|
||||
template <typename T> inline T GenerateRandomReal(const T min, const T max) {
|
||||
std::uniform_real_distribution<T> dist(min, max);
|
||||
return dist(GetRandomGenerator());
|
||||
}
|
||||
|
||||
inline int generate_width(int warp_size) {
|
||||
int exponent = 0;
|
||||
while (warp_size >>= 1) {
|
||||
++exponent;
|
||||
}
|
||||
|
||||
return GENERATE_COPY(map([](int e) { return 1 << e; }, range(1, exponent + 1)));
|
||||
}
|
||||
|
||||
inline uint64_t get_active_mask(unsigned int warp_id, unsigned int warp_size) {
|
||||
uint64_t active_mask = 0;
|
||||
switch (warp_id % 5) {
|
||||
case 0: // even threads in the warp
|
||||
active_mask = 0xAAAAAAAAAAAAAAAA;
|
||||
break;
|
||||
case 1: // odd threads in the warp
|
||||
active_mask = 0x5555555555555555;
|
||||
break;
|
||||
case 2: // first half of the warp
|
||||
for (int i = 0; i < warp_size / 2; i++) {
|
||||
active_mask = active_mask | (static_cast<uint64_t>(1) << i);
|
||||
}
|
||||
break;
|
||||
case 3: // second half of the warp
|
||||
for (int i = warp_size / 2; i < warp_size; i++) {
|
||||
active_mask = active_mask | (static_cast<uint64_t>(1) << i);
|
||||
}
|
||||
break;
|
||||
case 4: // all threads
|
||||
active_mask = 0xFFFFFFFFFFFFFFFF;
|
||||
break;
|
||||
}
|
||||
return active_mask;
|
||||
}
|
||||
|
||||
template <typename T, std::enable_if_t<std::is_integral<T>::value, bool> = true>
|
||||
inline T expandPrecision(int X) { return X; }
|
||||
|
||||
template <typename T, std::enable_if_t<std::is_floating_point<T>::value, bool> = true>
|
||||
inline T expandPrecision(int X) {
|
||||
return X * 3.141592653589793115997963468544185161590576171875;
|
||||
}
|
||||
|
||||
template <typename T, std::enable_if_t<std::is_same<T, __half>::value, bool> = true>
|
||||
inline __half expandPrecision(int X) {
|
||||
return (__half)expandPrecision<float>(X);
|
||||
}
|
||||
|
||||
template <typename T, std::enable_if_t<std::is_same<T, __half2>::value, bool> = true>
|
||||
inline __half2 expandPrecision(int X) {
|
||||
__half H = expandPrecision<float>(X);
|
||||
return {H, H};
|
||||
}
|
||||
|
||||
template <typename T, std::enable_if_t<std::is_integral<T>::value, bool> = true>
|
||||
inline void expandPrecision(T* Array, int size) {
|
||||
(void)Array;
|
||||
(void)size;
|
||||
}
|
||||
|
||||
template <typename T, std::enable_if_t<std::is_floating_point<T>::value, bool> = true>
|
||||
inline void expandPrecision(T *Array, int size) {
|
||||
for (int i = 0; i != size; ++i) {
|
||||
Array[i] *= 3.141592653589793115997963468544185161590576171875;
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
inline void initializeInput(T *Input, int size) {
|
||||
int Values[] = {0, -1, 2, 3, 4, 5, -6, 7,
|
||||
8, -9, 10, 11, 12, 13, -14, 15,
|
||||
16, 17, -18, 19, 20, -21, 22, 23,
|
||||
24, 25, 26, -27, 28, 29, 30, 31,
|
||||
-32, 33, 34, 35, -36, 37, 38, -39,
|
||||
40, 41, 42, 43, -44, -45, 46, 47,
|
||||
48, 49, 50, -51, 52, 53, -54, 55,
|
||||
56, 57, -58, 59, 60, 61, 62, -63};
|
||||
|
||||
for (int i = 0; i != size; ++i) {
|
||||
Input[i] = expandPrecision<T>(Values[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
inline void initializeExpected(T *Expected, int *Values, int size) {
|
||||
for (int i = 0; i != size; ++i) {
|
||||
Expected[i] = expandPrecision<T>(Values[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
inline bool compareEqual(T X, T Y) { return X == Y; }
|
||||
|
||||
template <>
|
||||
inline bool compareEqual(__half X, __half Y) {
|
||||
return __half2float(X) == __half2float(Y);
|
||||
}
|
||||
|
||||
template <>
|
||||
inline bool compareEqual(__half2 X, __half2 Y) {
|
||||
return compareEqual(X.x, Y.x) && compareEqual(X.y, Y.y);
|
||||
}
|
||||
|
||||
inline bool compareMaskEqual(unsigned long long *Actual, unsigned long long *Expected,
|
||||
int i, int warpSize) {
|
||||
if (warpSize == 32)
|
||||
return (unsigned)Actual[i] == (unsigned)Expected[i];
|
||||
return Actual[i] == Expected[i];
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
inline T alignUp(T num, size_t n) {
|
||||
if (num % n == 0) {
|
||||
return num;
|
||||
}
|
||||
|
||||
return ((num + n - 1) / n) * n;
|
||||
}
|
||||
@@ -0,0 +1,230 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
#define HIP_ENABLE_WARP_SYNC_BUILTINS
|
||||
#define HIP_ENABLE_EXTRA_WARP_SYNC_TYPES
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include "warp_common.hh"
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_fp16.h>
|
||||
#include <resource_guards.hh>
|
||||
#include <memory>
|
||||
#include <vector>
|
||||
#include <functional>
|
||||
#include <algorithm>
|
||||
#include <cstdlib>
|
||||
#include <cmd_options.hh>
|
||||
#include <tuple>
|
||||
|
||||
#define NELEMS(array) (sizeof(array) / sizeof(array[0]))
|
||||
|
||||
template <class T>
|
||||
// @input an array containing one value per lane to be used as input for the reduction
|
||||
// @masks a list of masks, none of them sharing bits
|
||||
__global__ void multipleMasksKernel(T* output, const T* input, const unsigned long long* masks, int numMasks)
|
||||
{
|
||||
bool isInAnyOfTheMasks = false;
|
||||
int numMask = 0;
|
||||
unsigned long long mask;
|
||||
|
||||
while (numMask < numMasks && !isInAnyOfTheMasks) {
|
||||
mask = masks[numMask];
|
||||
if ((1ul << threadIdx.x) & mask)
|
||||
isInAnyOfTheMasks = true;
|
||||
|
||||
numMask++;
|
||||
}
|
||||
|
||||
if (!isInAnyOfTheMasks)
|
||||
return;
|
||||
|
||||
output[threadIdx.x] = __reduce_add_sync<decltype(mask)>(mask, input[threadIdx.x]);
|
||||
}
|
||||
|
||||
template <class T, class Op, class MaskType>
|
||||
__global__ void reduceOp(T* output, const T* input, const MaskType* masks, int numReduces, Op)
|
||||
{
|
||||
int tid = threadIdx.x;
|
||||
|
||||
for (int i = 0; i < numReduces; i++) {
|
||||
if (masks[i] & (1ul << tid)) {
|
||||
// call the operator only if the lane is mentioned in the mask
|
||||
T& result = output[warpSize * i + tid];
|
||||
|
||||
if constexpr (std::is_same<Op, std::plus<T>>::value)
|
||||
result = __reduce_add_sync(masks[i], input[tid]);
|
||||
else if constexpr (std::is_same<Op, MinOp<T>>::value)
|
||||
result = __reduce_min_sync(masks[i], input[tid]);
|
||||
else if constexpr (std::is_same<Op, MaxOp<T>>::value)
|
||||
result = __reduce_max_sync(masks[i], input[tid]);
|
||||
else if constexpr (std::is_same<Op, std::logical_and<T>>::value)
|
||||
result = __reduce_and_sync(masks[i], input[tid]);
|
||||
else if (std::is_same<Op, std::logical_or<T>>::value)
|
||||
result = __reduce_or_sync(masks[i], input[tid]);
|
||||
else if (std::is_same<Op, XorOp<T>>::value)
|
||||
result = __reduce_xor_sync(masks[i], input[tid]);
|
||||
else
|
||||
assert(false && "Unsupported operator");
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
template <class T>
|
||||
void runTestMultipleMasks(unsigned long long masks[], int numMasks)
|
||||
{
|
||||
using namespace Catch::Matchers;
|
||||
using distribution = typename DistributionType<T>::type;
|
||||
unsigned int wavefrontSize = getWarpSize();
|
||||
LinearAllocGuard<unsigned long long> d_masks(LinearAllocs::hipMalloc, numMasks * sizeof(decltype(masks[0])));
|
||||
LinearAllocGuard<T> d_input, input;
|
||||
LinearAllocGuard<T> output(LinearAllocs::malloc, wavefrontSize * sizeof(T));
|
||||
LinearAllocGuard<T> d_output(LinearAllocs::hipMalloc, wavefrontSize * sizeof(T));
|
||||
std::plus<T> op;
|
||||
std::mt19937_64 gen(123);
|
||||
T a = std::is_same<T, half>::value? std::numeric_limits<unsigned short>::lowest() : -1023;
|
||||
T b = std::is_same<T, half>::value? std::numeric_limits<unsigned short>::max() : 1023;
|
||||
distribution distInput(a, b);
|
||||
dim3 blkDim { wavefrontSize };
|
||||
dim3 grdDim { 1u };
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_masks.ptr(), &masks[0], d_masks.size_bytes(), hipMemcpyHostToDevice));
|
||||
genRandomBuffers(d_input, input, distInput, gen, wavefrontSize);
|
||||
multipleMasksKernel<T><<<grdDim, blkDim>>>(d_output.ptr(), d_input.ptr(), d_masks.ptr(), numMasks);
|
||||
HIP_CHECK(hipMemcpy(output.ptr(), d_output.ptr(), d_output.size_bytes(), hipMemcpyDeviceToHost));
|
||||
|
||||
for (int numMask = 0; numMask < numMasks; numMask++) {
|
||||
unsigned long long mask = masks[numMask];
|
||||
T expected = calculateExpected<T>(input.ptr(), op, mask);
|
||||
int lane = 0;
|
||||
|
||||
while (lane < wavefrontSize) {
|
||||
if ((1ul << lane) & mask) {
|
||||
T result = output.ptr()[lane];
|
||||
|
||||
if constexpr (std::is_integral<T>::value) {
|
||||
// for integral types the result should match exactly
|
||||
if (result != expected) {
|
||||
printMismatch(result, expected, input.ptr(), mask);
|
||||
REQUIRE(result == expected);
|
||||
}
|
||||
} else
|
||||
compareFloatingPoint(result, expected, mask, input.ptr());
|
||||
}
|
||||
|
||||
lane++;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Unit_hipReduceSingleMasks", "", int, unsigned int, long long, unsigned long long, float, half, double) {
|
||||
unsigned long long fullMask = getWarpSize() == 64? ~0ul : 0xFFFFFFFF;
|
||||
unsigned long long oneBitMasks[] = { 0b1 & fullMask};
|
||||
unsigned long long everyFifthMasks[] = { Every5thBit & fullMask };
|
||||
unsigned long long everyNinethMasks[] = { Every9thBit & fullMask };
|
||||
unsigned long long everyFifthButNinethMasks[] = { Every5thBut9th & fullMask};
|
||||
|
||||
runTestMultipleMasks<TestType>(oneBitMasks, NELEMS(oneBitMasks));
|
||||
runTestMultipleMasks<TestType>(everyFifthMasks, NELEMS(everyFifthMasks));
|
||||
runTestMultipleMasks<TestType>(everyNinethMasks, NELEMS(everyNinethMasks));
|
||||
runTestMultipleMasks<TestType>(everyFifthButNinethMasks, NELEMS(everyFifthButNinethMasks));
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Unit_hipReduceMultipleMasks", "",
|
||||
int, unsigned int, long long, unsigned long long, float, half, double) {
|
||||
if (getWarpSize() == 64) {
|
||||
unsigned long long masks[] = { 0b0110011, 0x0F0F0F0F00000000, 0xF0F0F0F000000000,
|
||||
0x000000000F0F0F00, 0b0000100};
|
||||
// these divergent masks, when combined, occupy the whole set of lanes
|
||||
unsigned long long fullMasks[] = { 0xFFFF000000000000, 0x0000FFFFFFFF0000, 0x000000000000FFFF};
|
||||
unsigned long long fullMasksEvenOdd[] = { 0x5555555555555555, // even lanes
|
||||
0xAAAAAAAAAAAAAAAA }; // odd lanes
|
||||
|
||||
runTestMultipleMasks<TestType>(masks, NELEMS(masks));
|
||||
runTestMultipleMasks<TestType>(fullMasks, NELEMS(fullMasks));
|
||||
runTestMultipleMasks<TestType>(fullMasksEvenOdd, NELEMS(fullMasksEvenOdd));
|
||||
} else {
|
||||
unsigned long long masks1[] = { 0x0F0F0F0F, 0xF0F0F0F0 };
|
||||
unsigned long long masks2[] = { 0b0110011, 0x0F0F0F00, 0b0000100};
|
||||
runTestMultipleMasks<TestType>(masks1, NELEMS(masks1));
|
||||
runTestMultipleMasks<TestType>(masks2, NELEMS(masks2));
|
||||
}
|
||||
}
|
||||
|
||||
template <template <typename> class Op, class Type = void>
|
||||
void runTestReduceForTypes(const std::tuple<>)
|
||||
{
|
||||
}
|
||||
|
||||
template <template <typename> class Op, class T, typename... Types>
|
||||
void runTestReduceForTypes(const std::tuple<T, Types...>)
|
||||
{
|
||||
unsigned int wavefrontSize = getWarpSize();
|
||||
dim3 blkDim { wavefrontSize };
|
||||
dim3 grdDim { 1u };
|
||||
std::tuple<Types...> remainingTypes;
|
||||
int iteration = 0;
|
||||
auto reduceFunc = [&](T* d_output, const T* d_input, const unsigned long long* d_masks, int numReduces, Op<T> op) {
|
||||
reduceOp<T><<<grdDim, blkDim>>>(d_output, d_input, d_masks, numReduces, op);
|
||||
};
|
||||
bool customNumIterations = cmd_options.reduce_iterations != 1;
|
||||
|
||||
if (customNumIterations)
|
||||
std::cout << "\n" << opToString<T, Op>() << " - " << typeToString<T>() << "\n";
|
||||
|
||||
while (iteration < cmd_options.reduce_iterations) {
|
||||
runTestReduce<T, decltype(reduceFunc), Op>(iteration, reduceFunc);
|
||||
iteration++;
|
||||
|
||||
if (customNumIterations) {
|
||||
std::cout << "\rIteration: " << iteration;
|
||||
std::flush(std::cout);
|
||||
}
|
||||
}
|
||||
|
||||
runTestReduceForTypes<Op>(remainingTypes);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipReduceRandom") {
|
||||
const std::tuple<int, unsigned int, long long, unsigned long long, float, half, double> allTypes;
|
||||
const std::tuple<int, unsigned int, long long, unsigned long long> integralTypes;
|
||||
|
||||
SECTION("add") {
|
||||
runTestReduceForTypes<std::plus>(allTypes);
|
||||
}
|
||||
|
||||
SECTION("min") {
|
||||
runTestReduceForTypes<MinOp>(allTypes);
|
||||
}
|
||||
|
||||
SECTION("max") {
|
||||
runTestReduceForTypes<MaxOp>(allTypes);
|
||||
}
|
||||
|
||||
SECTION("and") {
|
||||
runTestReduceForTypes<std::logical_and>(integralTypes);
|
||||
}
|
||||
|
||||
SECTION("or") {
|
||||
runTestReduceForTypes<std::logical_or>(integralTypes);
|
||||
}
|
||||
|
||||
SECTION("xor") {
|
||||
runTestReduceForTypes<XorOp>(integralTypes);
|
||||
}
|
||||
}
|
||||
In neuem Issue referenzieren
Einen Benutzer sperren