From 97c2b31fda9fbb285f2f5cd4709a5ffb2bdceb61 Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Wed, 30 Sep 2026 18:35:38 +0000 Subject: [PATCH 1/7] fix(miopen): log reduce tensor lengths as size_t in LogCmdRedux The lambda that stringifies the input lengths for the logged MIOpenDriver command took an int, so lengths above INT32_MAX were printed truncated (e.g. 2^31 as -2147483648) and the logged command did not reproduce the failing call. Co-Authored-By: Claude Opus 5.5 (1M context) --- projects/miopen/src/reducetensor_api.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/miopen/src/reducetensor_api.cpp b/projects/miopen/src/reducetensor_api.cpp index 07621bd1b24a..569f23456567 100644 --- a/projects/miopen/src/reducetensor_api.cpp +++ b/projects/miopen/src/reducetensor_api.cpp @@ -61,7 +61,7 @@ static void LogCmdRedux(const miopen::ReduceTensorDescriptor reduceTensorDesc, std::transform(inLens.begin(), inLens.end(), std::back_inserter(strInLens), - [](const int x) { return std::to_string(x); }); + [](std::size_t x) { return std::to_string(x); }); ss << " -D " << miopen::JoinStrings(strInLens, ","); ss << " -I " << reduceTensorDesc.reduceTensorIndices_; ss << " -N " << reduceTensorDesc.reduceTensorNanOpt_; From 7a864bbe652dcd9290222bb19ef74d1ee15167a0 Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Wed, 30 Sep 2026 19:21:26 +0000 Subject: [PATCH 2/7] fix(miopen): reject reduce tensors that do not fit the int32 kernels (ALMIOPEN-2718) The legacy reduction kernels take lengths, strides, element counts and grid sizes as int and address memory through 32-bit byte offsets, but ReduceTensor narrowed the size_t values with static_cast, silently producing wrong results for large tensors. - ReduceTensor now rejects with miopenStatusBadParm any input or output tensor whose lengths or strides exceed INT32_MAX, or whose byte span does, plus element counts and grid sizes above INT32_MAX. Strides of length-1 dimensions are exempt: they never contribute to an offset, and the kernel is passed stride 1 for them. - GetWorkspaceSize does not do the int32 checks, since the RNN uses it to size reductions that may never run. - Descriptor validation (rank match, rank <= 6, output lengths) is shared by GetWorkspaceSize, GetIndicesSize and ReduceTensor, so the sizing queries now also reject rank > 6. - All errors thrown from reducetensor.cpp use miopenStatusBadParm instead of miopenStatusUnknownError. - Fix int overflow in get_padding_need (warpwise branch). - Remove the unused GetIndicesTypeSize(). Co-Authored-By: Claude Opus 5.5 (1M context) --- projects/miopen/src/reducetensor.cpp | 171 ++++++++++++++++----------- 1 file changed, 104 insertions(+), 67 deletions(-) diff --git a/projects/miopen/src/reducetensor.cpp b/projects/miopen/src/reducetensor.cpp index 3b4ee5e37477..62ec34928680 100644 --- a/projects/miopen/src/reducetensor.cpp +++ b/projects/miopen/src/reducetensor.cpp @@ -31,7 +31,9 @@ #include #include +#include #include +#include #include #include #include @@ -200,18 +202,6 @@ struct ReductionKernelConfigurator }; }; -inline int GetIndicesTypeSize(miopenIndicesType_t t) -{ - switch(t) - { - case MIOPEN_32BIT_INDICES: return (4); - case MIOPEN_64BIT_INDICES: return (8); - case MIOPEN_16BIT_INDICES: return (2); - case MIOPEN_8BIT_INDICES: return (1); - } - MIOPEN_THROW("Unknown data type"); -} - inline int GetDataTypeSize(miopenDataType_t t) { switch(t) @@ -228,7 +218,8 @@ inline int GetDataTypeSize(miopenDataType_t t) case miopenInt64: break; }; - MIOPEN_THROW("Only float, half, double, bfloat16, int8 data types are supported."); + MIOPEN_THROW(miopenStatusBadParm, + "Only float, half, double, bfloat16, int8 data types are supported."); }; }; // end of namespace detail @@ -253,7 +244,7 @@ static ck::DataTypeEnum_t mapDataTypeId(miopenDataType_t t) case miopenInt64: break; }; - MIOPEN_THROW("Only float, half, double data type is supported."); + MIOPEN_THROW(miopenStatusBadParm, "Only float, half, double data type is supported."); }; static ck::ReduceTensorOp_t mapReduceOpId(miopenReduceTensorOp_t t) @@ -272,7 +263,7 @@ static ck::ReduceTensorOp_t mapReduceOpId(miopenReduceTensorOp_t t) case MIOPEN_REDUCE_TENSOR_NORM2: return ReduceTensorOp_t::NORM2; }; - MIOPEN_THROW("Operation is not supported"); + MIOPEN_THROW(miopenStatusBadParm, "Operation is not supported"); }; static std::string get_network_config_string_from_type_enums(miopenDataType_t TSrc, @@ -387,8 +378,10 @@ static std::pair get_padding_need(ReductionMethod_t reduceImpl, case Reduce_DirectWarpWise: copySliceLen = warpSize * tunable->GredAccessesPerThreadInWarp; src_need_padding = - (invariantLen < GridSize * BlockSize / warpSize || toReduceLen % copySliceLen > 0); - dst_need_padding = (invariantLen < GridSize * BlockSize / warpSize); + (invariantLen < static_cast(GridSize) * BlockSize / warpSize || + toReduceLen % copySliceLen > 0); + dst_need_padding = + (invariantLen < static_cast(GridSize) * BlockSize / warpSize); return std::make_pair(src_need_padding, dst_need_padding); case Reduce_BlockWise: @@ -447,33 +440,93 @@ ReduceTensorDescriptor::ReduceTensorDescriptor(miopenReduceTensorOp_t reduceTens { if(reduceTensorIndices == MIOPEN_REDUCE_TENSOR_FLATTENED_INDICES && reduceTensorIndicesType != MIOPEN_32BIT_INDICES) - MIOPEN_THROW("Only int32 type is supported for ReduceTensor indices."); + MIOPEN_THROW(miopenStatusBadParm, "Only int32 type is supported for ReduceTensor indices."); }; -// This is WS requirement of the dynamic reduction. -// We must enforce it especially when reduction is used internally. -constexpr std::size_t workspaceAlignRequirementBytes = 64; +namespace { -// return the size of the workspace in bytes, so that the workspace buffer can be prepared by the -// user -std::size_t ReduceTensorDescriptor::GetWorkspaceSize(const Handle& handle, - const TensorDescriptor& inDesc, - const TensorDescriptor& outDesc) const +void ValidateReduceDescriptors(const TensorDescriptor& inDesc, const TensorDescriptor& outDesc) { + // the legacy reduction kernels take at most 6 dimensions + constexpr std::size_t maxReduceRank = 6; + const auto& inDescLengths = inDesc.GetLengths(); const auto& outDescLengths = outDesc.GetLengths(); if(inDescLengths.size() != outDescLengths.size()) - MIOPEN_THROW("The number of dimensions of the input and output tensor should match."); + MIOPEN_THROW(miopenStatusBadParm, + "The number of dimensions of the input and output tensor should match."); - for(int i = 0; i < inDescLengths.size(); i++) + if(inDescLengths.size() > maxReduceRank) + MIOPEN_THROW(miopenStatusBadParm, + "Invalid TensorDescriptor, at most number of dimensions of 6 is supported."); + + for(std::size_t i = 0; i < inDescLengths.size(); i++) { if(outDescLengths[i] != 1 && outDescLengths[i] != inDescLengths[i]) { - MIOPEN_THROW("The length of the output tensor dimension should either be 1 or be equal " + MIOPEN_THROW(miopenStatusBadParm, + "The length of the output tensor dimension should either be 1 or be equal " "to the length of the corresponding dimension of the input tensor."); } - }; + } +} + +// The kernels receive lengths and strides as int and address global memory through 32-bit byte +// offsets (legacy CK buffer addressing), so both have to fit into int. The stride of a length-1 +// dimension never contributes to an offset, so it is not restricted. +void ValidateReduceTensorFitsIntoInt(const TensorDescriptor& desc, const char* name) +{ + const auto& lengths = desc.GetLengths(); + const auto& strides = desc.GetStrides(); + + // offset of the last element; the terms are capped, so the sum cannot wrap (a capped term + // already makes the span exceed INT32_MAX bytes) + std::size_t lastOffset = 0; + for(std::size_t i = 0; i < lengths.size(); i++) + { + if(lengths[i] == 1) + continue; + + if(lengths[i] > std::numeric_limits::max() || + strides[i] > std::numeric_limits::max()) + MIOPEN_THROW(miopenStatusBadParm, + std::string{"Reduction "} + name + + " tensor has a length or stride exceeding INT32_MAX, which is not " + "supported."); + + lastOffset += + std::min((lengths[i] - 1) * strides[i], std::numeric_limits::max()); + } + + if((lastOffset + 1) * detail::GetDataTypeSize(desc.GetType()) > std::numeric_limits::max()) + MIOPEN_THROW(miopenStatusBadParm, + std::string{"Reduction "} + name + + " tensor spans more than INT32_MAX bytes, which is not supported."); +} + +int ToIntOrThrow(std::size_t value, const char* what) +{ + if(value > std::numeric_limits::max()) + MIOPEN_THROW(miopenStatusBadParm, + std::string{"Reduction "} + what + " " + std::to_string(value) + + " exceeds INT32_MAX, which is not supported."); + return static_cast(value); +} + +} // namespace + +// This is WS requirement of the dynamic reduction. +// We must enforce it especially when reduction is used internally. +constexpr std::size_t workspaceAlignRequirementBytes = 64; + +// return the size of the workspace in bytes, so that the workspace buffer can be prepared by the +// user +std::size_t ReduceTensorDescriptor::GetWorkspaceSize(const Handle& handle, + const TensorDescriptor& inDesc, + const TensorDescriptor& outDesc) const +{ + ValidateReduceDescriptors(inDesc, outDesc); auto invariantLength = outDesc.GetElementSize(); auto toReduceLength = inDesc.GetElementSize() / invariantLength; @@ -508,20 +561,7 @@ std::size_t ReduceTensorDescriptor::GetWorkspaceSize(const Handle& handle, std::size_t ReduceTensorDescriptor::GetIndicesSize(const TensorDescriptor& inDesc, const TensorDescriptor& outDesc) const { - const auto& inDescLengths = inDesc.GetLengths(); - const auto& outDescLengths = outDesc.GetLengths(); - - if(inDescLengths.size() != outDescLengths.size()) - MIOPEN_THROW("The number of dimensions of the input and output tensor should match."); - - for(int i = 0; i < inDescLengths.size(); i++) - { - if(outDescLengths[i] != 1 && outDescLengths[i] != inDescLengths[i]) - { - MIOPEN_THROW("The length of the output tensor dimension should either be 1 or be equal " - "to the length of the corresponding dimension of the input tensor."); - } - }; + ValidateReduceDescriptors(inDesc, outDesc); auto reduceIndicesOpt = this->reduceTensorIndices_; auto reduceOp = this->reduceTensorOp_; @@ -571,32 +611,24 @@ void ReduceTensorDescriptor::ReduceTensor(const Handle& handle, (reduceOp == MIOPEN_REDUCE_TENSOR_MIN || reduceOp == MIOPEN_REDUCE_TENSOR_MAX || reduceOp == MIOPEN_REDUCE_TENSOR_AMAX); - if(inDescLengths.size() > 6) - MIOPEN_THROW("Invalid TensorDescriptor, at most number of dimensions of 6 is supported."); + ValidateReduceDescriptors(aDesc, cDesc); if(need_indices && (reduceIndicesType != MIOPEN_32BIT_INDICES)) - MIOPEN_THROW("Only int32 type can be used for ReduceTensor indices."); - - if(inDescLengths.size() != outDescLengths.size()) - MIOPEN_THROW("The number of dimensions of the input and output tensor should match."); + MIOPEN_THROW(miopenStatusBadParm, "Only int32 type can be used for ReduceTensor indices."); - for(int i = 0; i < inDescLengths.size(); i++) - { - if(outDescLengths[i] != 1 && outDescLengths[i] != inDescLengths[i]) - { - MIOPEN_THROW("The length of the output tensor dimension should either be 1 or be equal " - "to the length of the corresponding dimension of the input tensor."); - } - }; + // GetWorkspaceSize() deliberately does not check this: it is also used just to size buffers + // for reductions that are never run (e.g. RNN inference) + ValidateReduceTensorFitsIntoInt(aDesc, "input"); + ValidateReduceTensorFitsIntoInt(cDesc, "output"); std::size_t ws_sizeInBytes = this->GetWorkspaceSize(handle, aDesc, cDesc); std::size_t indices_sizeInBytes = this->GetIndicesSize(aDesc, cDesc); if(ws_sizeInBytes > workspaceSizeInBytes) - MIOPEN_THROW("The workspace size allocated is not enough!"); + MIOPEN_THROW(miopenStatusBadParm, "The workspace size allocated is not enough!"); if(indices_sizeInBytes > indicesSizeInBytes) - MIOPEN_THROW("The indices size allocated is not enough!"); + MIOPEN_THROW(miopenStatusBadParm, "The indices size allocated is not enough!"); // invariantLength and toReduceLength are used to determine the kernel configuration const auto invariantLength = cDesc.GetElementSize(); @@ -614,7 +646,11 @@ void ReduceTensorDescriptor::ReduceTensor(const Handle& handle, const ReductionMethod_t reduceImpl = configurator.getReductionMethod(invariantLength, toReduceLength); - const int gridSize = configurator.getGridSize(invariantLength, toReduceLength); + // element counts are only bounded by the byte span checks above for non-overlapping tensors + ToIntOrThrow(invariantLength, "output element count"); + const int origReduceLen = ToIntOrThrow(toReduceLength, "reduced element count"); + const int gridSize = + ToIntOrThrow(configurator.getGridSize(invariantLength, toReduceLength), "grid size"); const int blkGroupSize = (reduceImpl == Reduce_MultiBlock) ? static_cast(gridSize / invariantLength) : 0; @@ -633,7 +669,8 @@ void ReduceTensorDescriptor::ReduceTensor(const Handle& handle, if(toReduceDims.empty()) { - MIOPEN_THROW("Invalid TensorDescriptor, at least one dimension of the input tensor should " + MIOPEN_THROW(miopenStatusBadParm, + "Invalid TensorDescriptor, at least one dimension of the input tensor should " "be reduced."); } @@ -647,13 +684,12 @@ void ReduceTensorDescriptor::ReduceTensor(const Handle& handle, : *reinterpret_cast(beta); { // use dynamic reduction - const int origReduceLen = toReduceLength; - int p_inLengths[6] = {0}; int p_inStrides[6] = {0}; int p_outLengths[6] = {0}; int p_outStrides[6] = {0}; + // the casts below are safe, see ValidateReduceTensorFitsIntoInt() int pos = 0; for(int i = 0; i < outDescLengths.size(); i++) { @@ -674,7 +710,8 @@ void ReduceTensorDescriptor::ReduceTensor(const Handle& handle, if(outDescLengths[i] == 1) { p_inLengths[pos] = static_cast(inDescLengths[i]); - p_inStrides[pos] = static_cast(inDescStrides[i]); + // a length-1 stride is unused and may not fit into int + p_inStrides[pos] = inDescLengths[i] == 1 ? 1 : static_cast(inDescStrides[i]); pos++; }; }; @@ -831,8 +868,8 @@ void ReduceTensorDescriptor::ReduceTensor(const Handle& handle, if(useTwoCalls) { const auto toReduceLength_2 = blkGroupSize; - const int gridSize_2 = - static_cast(configurator.getGridSize_2(invariantLength, toReduceLength_2)); + const int gridSize_2 = ToIntOrThrow( + configurator.getGridSize_2(invariantLength, toReduceLength_2), "grid size"); const std::vector vgd2_2 = { static_cast(gridSize_2) * tunable->BlockSize, size_t{1}, size_t{1}}; const auto reduceImpl2 = configurator.GetReductionMethod_2(toReduceLength_2); From 5f47078581882f8fa9a890ce9b0e0ff562b64955 Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Wed, 30 Sep 2026 19:53:18 +0000 Subject: [PATCH 3/7] test(miopen): add reduce shape limit regression tests (ALMIOPEN-2718) New Smoke/GPU_ReduceShapeLimits_FP32 suite in its own file. It checks that ReduceTensor rejects with miopenStatusBadParm tensors that do not fit the int32 legacy kernels (rank > 6, input/output lengths or strides above INT32_MAX, byte span above INT32_MAX), that GetWorkspaceSize and GetIndicesSize reject rank > 6 but stay permissive otherwise (RNN sizing path), and that a length-1 dimension with a stride above INT32_MAX still reduces correctly on the GPU. Co-Authored-By: Claude Opus 5.5 (1M context) --- .../miopen/test/gtest/reduce_shape_limits.cpp | 188 ++++++++++++++++++ 1 file changed, 188 insertions(+) create mode 100644 projects/miopen/test/gtest/reduce_shape_limits.cpp diff --git a/projects/miopen/test/gtest/reduce_shape_limits.cpp b/projects/miopen/test/gtest/reduce_shape_limits.cpp new file mode 100644 index 000000000000..b49053cdbc07 --- /dev/null +++ b/projects/miopen/test/gtest/reduce_shape_limits.cpp @@ -0,0 +1,188 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include +#include + +#include + +#include "get_handle.hpp" +#include "workspace.hpp" + +#include +#include +#include +#include +#include +#include + +// Regression tests for the reduce shape limits: shapes the legacy kernels cannot handle must be +// rejected with an error instead of being silently truncated to int32. The rejected shapes are +// only described, never allocated. +namespace { + +enum class ReduceCheck +{ + SizingRejects, // GetWorkspaceSize() and GetIndicesSize() throw + ReduceTensorRejects, // ReduceTensor() throws + SizingAccepts, // GetWorkspaceSize() and GetIndicesSize() do not throw + ReduceTensorAccepts, // ReduceTensor() runs and gives the correct result +}; + +struct ReduceShapeLimitsCase +{ + std::string name; + ReduceCheck check; + std::vector inLengths; + std::vector inStrides; + std::vector outLengths; + std::vector outStrides; + std::string message; + + friend std::ostream& operator<<(std::ostream& os, const ReduceShapeLimitsCase& c) + { + return os << c.name; + } +}; + +std::vector GetReduceShapeLimitsCases() +{ + constexpr std::size_t intMax = std::numeric_limits::max(); + const std::vector rank7(7, 2); + const std::vector rank7Strides{64, 32, 16, 8, 4, 2, 1}; + const std::vector rank7Out(7, 1); + const std::string rank7Message = "at most number of dimensions of 6"; + const std::string tooLarge = "input tensor has a length or stride exceeding INT32_MAX"; + + using enum ReduceCheck; + // clang-format off + return { + {"Rank7", SizingRejects, rank7, rank7Strides, rank7Out, rank7Out, rank7Message}, + {"Rank7ReduceTensor", ReduceTensorRejects, rank7, rank7Strides, rank7Out, rank7Out, rank7Message}, + {"LengthAboveIntMax", ReduceTensorRejects, {intMax + 1, 1}, {1, 1}, {1, 1}, {1, 1}, tooLarge}, + {"StrideAboveIntMax", ReduceTensorRejects, {2, 4}, {intMax + 1, 1}, {1, 4}, {4, 1}, tooLarge}, + {"OutputStrideAboveIntMax", ReduceTensorRejects, {2, 4}, {4, 1}, {2, 1}, {intMax + 1, 1}, + "output tensor has a length or stride exceeding INT32_MAX"}, + // every length and stride fits, but the fp32 byte span does not + {"ByteSpanAboveIntMax", ReduceTensorRejects, {2, (1ULL << 28) + 1}, {(1ULL << 28) + 1, 1}, {2, 1}, {1, 1}, + "input tensor spans more than INT32_MAX bytes"}, + // sizing stays permissive: the RNN sizes reductions it may never run (ReductionWorkspaceSize) + {"StrideAboveIntMaxSizing", SizingAccepts, {2, 4}, {intMax + 1, 1}, {1, 4}, {4, 1}, ""}, + {"RnnLikeSizing", SizingAccepts, {1, 8, 16}, {1ULL << 32, 16, 1}, {1, 1, 16}, {16, 16, 1}, ""}, + // the stride of a length-1 dimension is never used, so it may exceed INT32_MAX + {"RnnLikeReduceTensor", ReduceTensorAccepts, {1, 8, 16}, {1ULL << 32, 16, 1}, {1, 1, 16}, {16, 16, 1}, ""}, + }; + // clang-format on +} + +} // anonymous namespace + +class GPU_ReduceShapeLimits_FP32 : public ::testing::TestWithParam +{ +}; + +TEST_P(GPU_ReduceShapeLimits_FP32, EnforcesLimits) +{ + const auto& c = GetParam(); + auto&& handle = get_handle(); + + // MAX with flattened indices, so that GetIndicesSize() has something to size + const miopen::ReduceTensorDescriptor reduceDesc{MIOPEN_REDUCE_TENSOR_MAX, + miopenFloat, + MIOPEN_NOT_PROPAGATE_NAN, + MIOPEN_REDUCE_TENSOR_FLATTENED_INDICES, + MIOPEN_32BIT_INDICES}; + const miopen::TensorDescriptor inDesc{miopenFloat, c.inLengths, c.inStrides}; + const miopen::TensorDescriptor outDesc{miopenFloat, c.outLengths, c.outStrides}; + + const auto expectBadParm = [&](auto&& call) { + try + { + call(); + ADD_FAILURE() << "no exception thrown"; + } + catch(const miopen::Exception& e) + { + EXPECT_EQ(e.status, miopenStatusBadParm); + EXPECT_NE(e.message.find(c.message), std::string::npos) << e.message; + } + }; + + switch(c.check) + { + case ReduceCheck::SizingRejects: + expectBadParm([&] { std::ignore = reduceDesc.GetWorkspaceSize(handle, inDesc, outDesc); }); + expectBadParm([&] { std::ignore = reduceDesc.GetIndicesSize(inDesc, outDesc); }); + break; + case ReduceCheck::ReduceTensorRejects: { + const float alpha = 1.0f; + const float beta = 0.0f; + // oversized buffer sizes, so the "not enough" checks cannot fire first + expectBadParm([&] { + reduceDesc.ReduceTensor(handle, + nullptr, + std::numeric_limits::max(), + nullptr, + std::numeric_limits::max(), + &alpha, + inDesc, + nullptr, + &beta, + outDesc, + nullptr); + }); + break; + } + case ReduceCheck::SizingAccepts: + EXPECT_NO_THROW(std::ignore = reduceDesc.GetWorkspaceSize(handle, inDesc, outDesc)); + EXPECT_NO_THROW(std::ignore = reduceDesc.GetIndicesSize(inDesc, outDesc)); + break; + case ReduceCheck::ReduceTensorAccepts: { + // the input is 3-D with a leading length-1 dimension; only the elements reachable through + // the other two dimensions are allocated + const std::size_t rows = c.inLengths[1]; + const std::size_t cols = c.inLengths[2]; + std::vector input(rows * cols); + std::iota(input.begin(), input.end(), 0.0f); + + const miopen::ReduceTensorDescriptor addDesc{MIOPEN_REDUCE_TENSOR_ADD, + miopenFloat, + MIOPEN_NOT_PROPAGATE_NAN, + MIOPEN_REDUCE_TENSOR_NO_INDICES, + MIOPEN_32BIT_INDICES}; + auto input_dev = handle.Write(input); + auto output_dev = handle.Write(std::vector(cols, 0.0f)); + Workspace wspace{addDesc.GetWorkspaceSize(handle, inDesc, outDesc)}; + + const float alpha = 1.0f; + const float beta = 0.0f; + addDesc.ReduceTensor(handle, + nullptr, + 0, + wspace.ptr(), + wspace.size(), + &alpha, + inDesc, + input_dev.get(), + &beta, + outDesc, + output_dev.get()); + + const auto output = handle.Read(output_dev, cols); + for(std::size_t k = 0; k < cols; k++) + { + float expected = 0.0f; + for(std::size_t j = 0; j < rows; j++) + expected += input[j * cols + k]; + EXPECT_EQ(output[k], expected) << "column " << k; + } + break; + } + } +} + +INSTANTIATE_TEST_SUITE_P(Smoke, + GPU_ReduceShapeLimits_FP32, + ::testing::ValuesIn(GetReduceShapeLimitsCases()), + [](const auto& info) { return info.param.name; }); From 1e6c060fe3faf09da2c74833f41fb49cfcffd76c Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Wed, 30 Sep 2026 19:56:20 +0000 Subject: [PATCH 4/7] docs(miopen): add CHANGELOG entries for reduce int32 limits (ALMIOPEN-2718) Co-Authored-By: Claude Opus 5.5 (1M context) --- projects/miopen/CHANGELOG.md | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/projects/miopen/CHANGELOG.md b/projects/miopen/CHANGELOG.md index b309592784e9..790f80da9566 100644 --- a/projects/miopen/CHANGELOG.md +++ b/projects/miopen/CHANGELOG.md @@ -5,6 +5,16 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj ## MIOpen 3.6.2 for ROCm 10.2.0 +### Resolved Issues +* [Reduce] Fixed `miopenReduceTensor` silently truncating tensor lengths, strides, element counts + and byte spans above INT32_MAX, which produced incorrect results. Such tensors are now rejected + with `miopenStatusBadParm`. +* [Reduce] `miopenGetReductionWorkspaceSize` and `miopenGetReductionIndicesSize` now reject tensors + of rank greater than 6, matching `miopenReduceTensor`. +* [Reduce] Invalid arguments to the reduction API (mismatched descriptors, unsupported types, + undersized workspace or indices buffers) now return `miopenStatusBadParm` instead of + `miopenStatusUnknownError`. + ### Added * [Conv] Added gfx950 depthwise backward-weights (fp16/bf16) and gfx1250 depthwise kernels to the `ConvHipConv` solver (hipconv v0.3.1). From f7ceb1629166864d908f96297a0742b37a1390ae Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Thu, 1 Oct 2026 17:11:46 +0000 Subject: [PATCH 5/7] fix(miopen): reject reductions whose padded length does not fit int (ALMIOPEN-2718) For a multi-block reduction with a single output the number of blocks is not capped, so the padded reduced length the kernel computes as int (reduceSizePerBlock * BlkGroupSize, about twice the reduced length) can exceed INT32_MAX even when the tensor itself fits: e.g. a packed fp16 tensor of 2^30 - 1 elements reduced to one value. The kernel then reduces nothing and silently returns a wrong result. get_padding_need now rejects such reductions with miopenStatusBadParm. ToIntOrThrow moves to an unnamed namespace ahead of detailDynamic so get_padding_need can use it. Co-Authored-By: Claude Opus 5.5 (1M context) --- projects/miopen/src/reducetensor.cpp | 30 ++++++++++++++++++---------- 1 file changed, 19 insertions(+), 11 deletions(-) diff --git a/projects/miopen/src/reducetensor.cpp b/projects/miopen/src/reducetensor.cpp index 62ec34928680..f785cb37ac34 100644 --- a/projects/miopen/src/reducetensor.cpp +++ b/projects/miopen/src/reducetensor.cpp @@ -224,6 +224,19 @@ inline int GetDataTypeSize(miopenDataType_t t) }; // end of namespace detail +namespace { + +int ToIntOrThrow(std::size_t value, const char* what) +{ + if(value > std::numeric_limits::max()) + MIOPEN_THROW(miopenStatusBadParm, + std::string{"Reduction "} + what + " " + std::to_string(value) + + " exceeds INT32_MAX, which is not supported."); + return static_cast(value); +} + +} // namespace + namespace detailDynamic { static ck::DataTypeEnum_t mapDataTypeId(miopenDataType_t t) @@ -365,6 +378,7 @@ static std::pair get_padding_need(ReductionMethod_t reduceImpl, bool dst_need_padding = false; int copySliceLen; int reduceSizePerBlock; + size_t paddedReduceLen; switch(reduceImpl) { @@ -395,8 +409,11 @@ static std::pair get_padding_need(ReductionMethod_t reduceImpl, (((toReduceLen + BlkGroupSize - 1) / BlkGroupSize + copySliceLen - 1) / copySliceLen) * copySliceLen; - src_need_padding = - (toReduceLen < static_cast(reduceSizePerBlock) * BlkGroupSize); + // the kernel computes the padded length as int; with a single output the number of + // blocks is not capped, so it can exceed INT32_MAX even when the tensor itself fits + paddedReduceLen = static_cast(reduceSizePerBlock) * BlkGroupSize; + ToIntOrThrow(paddedReduceLen, "padded reduced length"); + src_need_padding = (toReduceLen < paddedReduceLen); return std::make_pair(src_need_padding, dst_need_padding); }; @@ -505,15 +522,6 @@ void ValidateReduceTensorFitsIntoInt(const TensorDescriptor& desc, const char* n " tensor spans more than INT32_MAX bytes, which is not supported."); } -int ToIntOrThrow(std::size_t value, const char* what) -{ - if(value > std::numeric_limits::max()) - MIOPEN_THROW(miopenStatusBadParm, - std::string{"Reduction "} + what + " " + std::to_string(value) + - " exceeds INT32_MAX, which is not supported."); - return static_cast(value); -} - } // namespace // This is WS requirement of the dynamic reduction. From 6e4a670b6f356e4206c023b1e8cce9f9cb1e83b6 Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Thu, 1 Oct 2026 17:14:15 +0000 Subject: [PATCH 6/7] docs(miopen): update CHANGELOG entries for reduce int32 limits (ALMIOPEN-2718) Mention single reductions of about 2^30 or more fp16, bf16 or int8 elements under Resolved Issues, and move the rank > 6 rejection in the sizing queries and the UnknownError -> BadParm status change to a new Changed section, since they are behavior changes rather than fixes. Co-Authored-By: Claude Opus 5.5 (1M context) --- projects/miopen/CHANGELOG.md | 15 +++++++++------ 1 file changed, 9 insertions(+), 6 deletions(-) diff --git a/projects/miopen/CHANGELOG.md b/projects/miopen/CHANGELOG.md index 790f80da9566..02d3cb179b91 100644 --- a/projects/miopen/CHANGELOG.md +++ b/projects/miopen/CHANGELOG.md @@ -7,18 +7,21 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj ### Resolved Issues * [Reduce] Fixed `miopenReduceTensor` silently truncating tensor lengths, strides, element counts - and byte spans above INT32_MAX, which produced incorrect results. Such tensors are now rejected - with `miopenStatusBadParm`. + and byte spans above INT32_MAX, and single reductions of about 2^30 or more fp16, bf16 or int8 + elements, which produced incorrect results. Such tensors are now rejected with + `miopenStatusBadParm`. + +### Added +* [Conv] Added gfx950 depthwise backward-weights (fp16/bf16) and gfx1250 depthwise + kernels to the `ConvHipConv` solver (hipconv v0.3.1). + +### Changed * [Reduce] `miopenGetReductionWorkspaceSize` and `miopenGetReductionIndicesSize` now reject tensors of rank greater than 6, matching `miopenReduceTensor`. * [Reduce] Invalid arguments to the reduction API (mismatched descriptors, unsupported types, undersized workspace or indices buffers) now return `miopenStatusBadParm` instead of `miopenStatusUnknownError`. -### Added -* [Conv] Added gfx950 depthwise backward-weights (fp16/bf16) and gfx1250 depthwise - kernels to the `ConvHipConv` solver (hipconv v0.3.1). - ### Removed * [Conv] Removed gfx803 convolution solver `ConvBinWinogradRxSFused` and its kernel sources. * [Conv] Removed the gfx803 code paths and kernel sources from `ConvBinWinograd3x3U` and `ConvBinWinogradRxS`. From b8a606d855f642f783baa8e0b20c724c0bd3772d Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Thu, 1 Oct 2026 17:16:48 +0000 Subject: [PATCH 7/7] test(miopen): cover the single-output padded length limit in reduce tests (ALMIOPEN-2718) Add a Smoke/GPU_ReduceShapeLimits_FP16 suite with a packed fp16 tensor of INT32_MAX / 2 elements reduced to one value: the tensor fits into INT32_MAX bytes, but the padded length of the multi-block kernel does not. The test body is shared between the FP32 and FP16 suites. Also replace the magic numbers with named limits (the fp32 byte span case now uses the smallest rejected row length) and clarify the comments on the sizing cases and on the null workspace that keeps a regression from launching kernels. Co-Authored-By: Claude Opus 5.5 (1M context) --- .../miopen/test/gtest/reduce_shape_limits.cpp | 75 +++++++++++++++---- 1 file changed, 59 insertions(+), 16 deletions(-) diff --git a/projects/miopen/test/gtest/reduce_shape_limits.cpp b/projects/miopen/test/gtest/reduce_shape_limits.cpp index b49053cdbc07..4d32bd534749 100644 --- a/projects/miopen/test/gtest/reduce_shape_limits.cpp +++ b/projects/miopen/test/gtest/reduce_shape_limits.cpp @@ -46,9 +46,11 @@ struct ReduceShapeLimitsCase } }; -std::vector GetReduceShapeLimitsCases() +std::vector GetReduceShapeLimitsFP32Cases() { constexpr std::size_t intMax = std::numeric_limits::max(); + // the smallest row length for which two fp32 rows span more than INT32_MAX bytes + constexpr std::size_t fp32SpanRowLen = (intMax / sizeof(float) + 1) / 2; const std::vector rank7(7, 2); const std::vector rank7Strides{64, 32, 16, 8, 4, 2, 1}; const std::vector rank7Out(7, 1); @@ -65,26 +67,39 @@ std::vector GetReduceShapeLimitsCases() {"OutputStrideAboveIntMax", ReduceTensorRejects, {2, 4}, {4, 1}, {2, 1}, {intMax + 1, 1}, "output tensor has a length or stride exceeding INT32_MAX"}, // every length and stride fits, but the fp32 byte span does not - {"ByteSpanAboveIntMax", ReduceTensorRejects, {2, (1ULL << 28) + 1}, {(1ULL << 28) + 1, 1}, {2, 1}, {1, 1}, + {"ByteSpanAboveIntMax", ReduceTensorRejects, {2, fp32SpanRowLen}, {fp32SpanRowLen, 1}, {2, 1}, {1, 1}, "input tensor spans more than INT32_MAX bytes"}, - // sizing stays permissive: the RNN sizes reductions it may never run (ReductionWorkspaceSize) + // the sizing queries do not check the int32 limits, because the RNN uses them to size + // reductions it may never run (ReductionWorkspaceSize) {"StrideAboveIntMaxSizing", SizingAccepts, {2, 4}, {intMax + 1, 1}, {1, 4}, {4, 1}, ""}, - {"RnnLikeSizing", SizingAccepts, {1, 8, 16}, {1ULL << 32, 16, 1}, {1, 1, 16}, {16, 16, 1}, ""}, - // the stride of a length-1 dimension is never used, so it may exceed INT32_MAX - {"RnnLikeReduceTensor", ReduceTensorAccepts, {1, 8, 16}, {1ULL << 32, 16, 1}, {1, 1, 16}, {16, 16, 1}, ""}, + {"RnnLikeSizing", SizingAccepts, {1, 8, 16}, {intMax + 1, 16, 1}, {1, 1, 16}, {16, 16, 1}, ""}, + // the stride of a length-1 dimension is never used, so it may exceed INT32_MAX without the + // tensor being rejected + {"RnnLikeReduceTensor", ReduceTensorAccepts, {1, 8, 16}, {intMax + 1, 16, 1}, {1, 1, 16}, {16, 16, 1}, ""}, }; // clang-format on } -} // anonymous namespace - -class GPU_ReduceShapeLimits_FP32 : public ::testing::TestWithParam +std::vector GetReduceShapeLimitsFP16Cases() { -}; + // the most 2-byte elements that fit into INT32_MAX bytes + constexpr std::size_t maxLen = std::numeric_limits::max() / 2; -TEST_P(GPU_ReduceShapeLimits_FP32, EnforcesLimits) + // a single reduction of that many elements is split over so many blocks that the padded length + // the kernel computes exceeds INT32_MAX, although the tensor itself fits + return { + {"SingleOutputPaddedLengthAboveIntMax", + ReduceCheck::ReduceTensorRejects, + {maxLen}, + {1}, + {1}, + {1}, + "padded reduced length"}, + }; +} + +void RunReduceShapeLimitsCase(const ReduceShapeLimitsCase& c, miopenDataType_t type) { - const auto& c = GetParam(); auto&& handle = get_handle(); // MAX with flattened indices, so that GetIndicesSize() has something to size @@ -93,8 +108,8 @@ TEST_P(GPU_ReduceShapeLimits_FP32, EnforcesLimits) MIOPEN_NOT_PROPAGATE_NAN, MIOPEN_REDUCE_TENSOR_FLATTENED_INDICES, MIOPEN_32BIT_INDICES}; - const miopen::TensorDescriptor inDesc{miopenFloat, c.inLengths, c.inStrides}; - const miopen::TensorDescriptor outDesc{miopenFloat, c.outLengths, c.outStrides}; + const miopen::TensorDescriptor inDesc{type, c.inLengths, c.inStrides}; + const miopen::TensorDescriptor outDesc{type, c.outLengths, c.outStrides}; const auto expectBadParm = [&](auto&& call) { try @@ -118,7 +133,9 @@ TEST_P(GPU_ReduceShapeLimits_FP32, EnforcesLimits) case ReduceCheck::ReduceTensorRejects: { const float alpha = 1.0f; const float beta = 0.0f; - // oversized buffer sizes, so the "not enough" checks cannot fire first + // oversized buffer sizes, so the "not enough" checks cannot fire first; the null + // workspace makes ReduceTensor() throw before launching any kernel if the checks are + // missing, so a regression fails the test instead of running on invalid buffers expectBadParm([&] { reduceDesc.ReduceTensor(handle, nullptr, @@ -139,6 +156,7 @@ TEST_P(GPU_ReduceShapeLimits_FP32, EnforcesLimits) EXPECT_NO_THROW(std::ignore = reduceDesc.GetIndicesSize(inDesc, outDesc)); break; case ReduceCheck::ReduceTensorAccepts: { + ASSERT_EQ(type, miopenFloat) << "only implemented for fp32"; // the input is 3-D with a leading length-1 dimension; only the elements reachable through // the other two dimensions are allocated const std::size_t rows = c.inLengths[1]; @@ -182,7 +200,32 @@ TEST_P(GPU_ReduceShapeLimits_FP32, EnforcesLimits) } } +} // anonymous namespace + +class GPU_ReduceShapeLimits_FP32 : public ::testing::TestWithParam +{ +}; + +class GPU_ReduceShapeLimits_FP16 : public ::testing::TestWithParam +{ +}; + +TEST_P(GPU_ReduceShapeLimits_FP32, EnforcesLimits) +{ + RunReduceShapeLimitsCase(GetParam(), miopenFloat); +} + +TEST_P(GPU_ReduceShapeLimits_FP16, EnforcesLimits) +{ + RunReduceShapeLimitsCase(GetParam(), miopenHalf); +} + INSTANTIATE_TEST_SUITE_P(Smoke, GPU_ReduceShapeLimits_FP32, - ::testing::ValuesIn(GetReduceShapeLimitsCases()), + ::testing::ValuesIn(GetReduceShapeLimitsFP32Cases()), + [](const auto& info) { return info.param.name; }); + +INSTANTIATE_TEST_SUITE_P(Smoke, + GPU_ReduceShapeLimits_FP16, + ::testing::ValuesIn(GetReduceShapeLimitsFP16Cases()), [](const auto& info) { return info.param.name; });