SWDEV-420237 - Add tests for reduce sync operations (#102)
[ROCm/hip-tests commit: 001affc6ae]
Este cometimento está contido em:
Hernandez, Gerardo
cometido por
GitHub
GitHub
ascendente
105f388d54
cometimento
5d5fa6fc2a
@@ -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);
|
||||
}
|
||||
}
|
||||
Criar uma nova questão referindo esta
Bloquear um utilizador