From 00a4350f2c6b3e8db66bbbc6242dba781f1cfdd0 Mon Sep 17 00:00:00 2001 From: CMiservaAMD Date: Thu, 29 Oct 2026 16:42:44 -0400 Subject: [PATCH 1/3] (CMCL) hipDNN data SDK: allow zero (broadcast) strides in Tensor A zero stride on a dimension of size greater than 1 is a broadcast: every index along that dimension reads the same memory. Tensor now accepts non-negative strides, sizes its allocation from the largest reachable offset, and reports a tensor with a broadcast dimension as not packed. isTensorPacked applies the same rule. Negative strides remain unsupported. ALMIOPEN-3109 Co-Authored-By: Claude Opus 5.5 (1M context) --- .../utilities/ShapeUtilities.hpp | 10 +++ .../hipdnn_data_sdk/utilities/Tensor.hpp | 26 +++++- .../tests/utilities/TestShapeUtilities.cpp | 17 ++++ .../data_sdk/tests/utilities/TestTensor.cpp | 84 ++++++++++++++++++- 4 files changed, 133 insertions(+), 4 deletions(-) diff --git a/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/ShapeUtilities.hpp b/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/ShapeUtilities.hpp index a8344e57de59..fcae016dde1f 100644 --- a/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/ShapeUtilities.hpp +++ b/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/ShapeUtilities.hpp @@ -306,6 +306,16 @@ inline bool isTensorPacked(const std::vector& dims, const std::vector 1 maps several indices to one element, and the + // span check below can still pass when other axes leave matching gaps. + for(size_t axis = 0; axis < dims.size(); ++axis) + { + if(dims[axis] > 1 && strides[axis] == 0) + { + return false; + } + } + // Calculate total element count const auto count = std::accumulate(dims.begin(), dims.end(), static_cast(1), std::multiplies<>()); diff --git a/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/Tensor.hpp b/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/Tensor.hpp index df3bfd4469d9..0f2b1f19dd19 100644 --- a/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/Tensor.hpp +++ b/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/Tensor.hpp @@ -837,6 +837,15 @@ class TensorBase : public ITensor bool computeIsPacked(const std::vector& dims, const std::vector& strides) const { + // A zero stride on an axis of extent > 1 maps several indices to one element. + // Its element space can still equal the item count, so check the strides first. + for(size_t axis = 0; axis < dims.size(); ++axis) + { + if(dims[axis] > 1 && strides[axis] == 0) + { + return false; + } + } // Item count = largest stride * item count in that dimension return (calculateItemCount(dims) == calculateElementSpace(dims, strides)); } @@ -878,7 +887,7 @@ class Tensor : public TensorBase { validateDimsAndStridesSameSize(); validateAllPositive(_dims, "dimension"); - validateAllPositive(_strides, "stride"); + validateAllNonNegative(_strides, "stride"); // Set packed flag after validations since it can be incorrect if dims/strides are invalid. _packed = TensorBase::computeIsPacked(dims, strides); @@ -993,6 +1002,21 @@ class Tensor : public TensorBase } } + void validateAllNonNegative(const std::vector& values, + const std::string& valueName) const + { + for(size_t i = 0; i < values.size(); ++i) + { + if(values[i] < 0) + { + std::ostringstream oss; + oss << "All " << valueName << "s must be non-negative. " << valueName << " " << i + << " is " << values[i]; + throw std::invalid_argument(oss.str()); + } + } + } + utilities::MigratableMemory _memory; std::vector _dims; std::vector _strides; diff --git a/projects/hipdnn/data_sdk/tests/utilities/TestShapeUtilities.cpp b/projects/hipdnn/data_sdk/tests/utilities/TestShapeUtilities.cpp index 1d0a4328cc35..b3711c0bd304 100644 --- a/projects/hipdnn/data_sdk/tests/utilities/TestShapeUtilities.cpp +++ b/projects/hipdnn/data_sdk/tests/utilities/TestShapeUtilities.cpp @@ -777,6 +777,23 @@ TEST(TestShapeUtils, IsTensorPackedOverlappingStrides) EXPECT_FALSE(isTensorPacked(dims, overlappingStrides)); } +TEST(TestShapeUtils, IsTensorPackedZeroStrideWithMatchingSpan) +{ + // count = 4, span = (2-1)*3 + (2-1)*0 + 1 = 4 + EXPECT_FALSE(isTensorPacked({2, 2}, {3, 0})); +} + +TEST(TestShapeUtils, IsTensorPackedZeroStrideBroadcast) +{ + EXPECT_FALSE(isTensorPacked({2, 3, 4}, {4, 0, 1})); + EXPECT_FALSE(isTensorPacked({2, 3}, {0, 0})); +} + +TEST(TestShapeUtils, IsTensorPackedZeroStrideOnUnitExtentAxis) +{ + EXPECT_TRUE(isTensorPacked({2, 1, 3}, {3, 0, 1})); +} + TEST(TestShapeUtils, IsTensorPackedAllOnesShape) { // Shape with all 1s should be packed diff --git a/projects/hipdnn/data_sdk/tests/utilities/TestTensor.cpp b/projects/hipdnn/data_sdk/tests/utilities/TestTensor.cpp index ea04feef55aa..465f05a4e1b6 100644 --- a/projects/hipdnn/data_sdk/tests/utilities/TestTensor.cpp +++ b/projects/hipdnn/data_sdk/tests/utilities/TestTensor.cpp @@ -790,12 +790,90 @@ TEST(TestTensor, ConstructorThrowsOnNegativeStride) EXPECT_THROW(const Tensor tensor(dims, strides), std::invalid_argument); } -TEST(TestTensor, ConstructorThrowsOnZeroStride) +TEST(TestTensor, ZeroStrideAxisReadsOneElement) { const std::vector dims = {2, 3, 4}; - const std::vector strides = {12, 0, 1}; + const std::vector strides = {4, 0, 1}; - EXPECT_THROW(const Tensor tensor(dims, strides), std::invalid_argument); + Tensor tensor(dims, strides); + + EXPECT_EQ(tensor.elementCount(), 24); + EXPECT_EQ(tensor.elementSpace(), 8); + EXPECT_FALSE(tensor.isPacked()); + EXPECT_EQ(tensor.getIndex(1, 0, 3), 7); + EXPECT_EQ(tensor.getIndex(1, 2, 3), 7); + + tensor.setHostValue(5.0F, 1, 0, 2); + for(int64_t i1 = 0; i1 < dims[1]; ++i1) + { + EXPECT_EQ(tensor.getHostValue(1, i1, 2), 5.0F); + } +} + +TEST(TestTensor, AllZeroStridesUseOneElement) +{ + Tensor tensor({2, 3}, {0, 0}); + + EXPECT_EQ(tensor.elementCount(), 6); + EXPECT_EQ(tensor.elementSpace(), 1); + EXPECT_FALSE(tensor.isPacked()); + + tensor.fillWithValue(3.0F); + EXPECT_EQ(tensor.getHostValue(1, 2), 3.0F); +} + +TEST(TestTensor, ZeroStrideIsNotPackedWhenElementSpaceEqualsItemCount) +{ + const Tensor tensor({2, 2}, {3, 0}); + + ASSERT_EQ(tensor.elementCount(), tensor.elementSpace()); + EXPECT_FALSE(tensor.isPacked()); +} + +TEST(TestTensor, ZeroStrideOnUnitExtentAxisStaysPacked) +{ + const Tensor tensor({2, 1, 3}, {3, 0, 1}); + + EXPECT_TRUE(tensor.isPacked()); +} + +TEST(TestTensor, FillWithRandomValuesBroadcastsAlongZeroStrideAxis) +{ + const std::vector dims = {3, 4, 2}; + Tensor tensor(dims, {2, 0, 1}); + + tensor.fillWithRandomValues(-1.0F, 1.0F, 1234); + + for(int64_t i0 = 0; i0 < dims[0]; ++i0) + { + for(int64_t i2 = 0; i2 < dims[2]; ++i2) + { + const float first = tensor.getHostValue(i0, 0, i2); + EXPECT_GE(first, -1.0F); + EXPECT_LE(first, 1.0F); + for(int64_t i1 = 1; i1 < dims[1]; ++i1) + { + EXPECT_EQ(tensor.getHostValue(i0, i1, i2), first); + } + } + } +} + +TEST(TestTensor, ZeroStrideTensorRoundTripsThroughDevice) +{ + SKIP_IF_NO_DEVICES(); + + Tensor tensor({2, 3}, {1, 0}); + tensor.setHostValue(1.5F, 0, 0); + tensor.setHostValue(2.5F, 1, 0); + + auto* device = tensor.rawDeviceData(); + std::vector fromDevice(tensor.elementSpace()); + ASSERT_EQ( + hipMemcpy( + fromDevice.data(), device, fromDevice.size() * sizeof(float), hipMemcpyDeviceToHost), + hipSuccess); + EXPECT_EQ(fromDevice, (std::vector{1.5F, 2.5F})); } TEST(TestTensor, ConstructorThrowsOnMultipleInvalidDimensions) From 6b7c38da71b59d53e92abf7dd6095728e99c053a Mon Sep 17 00:00:00 2001 From: CMiservaAMD Date: Thu, 29 Oct 2026 17:19:40 -0400 Subject: [PATCH 2/3] (CMCL) hipDNN data SDK: share the broadcast-axis check and tighten zero-stride tests ALMIOPEN-3109 Co-Authored-By: Claude Opus 5.5 (1M context) --- .../utilities/ShapeUtilities.hpp | 23 ++++++--- .../hipdnn_data_sdk/utilities/Tensor.hpp | 47 ++++++------------- .../data_sdk/tests/utilities/TestTensor.cpp | 21 ++++++--- 3 files changed, 46 insertions(+), 45 deletions(-) diff --git a/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/ShapeUtilities.hpp b/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/ShapeUtilities.hpp index fcae016dde1f..3fb168cd50a7 100644 --- a/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/ShapeUtilities.hpp +++ b/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/ShapeUtilities.hpp @@ -290,6 +290,20 @@ inline std::vector return generateStrides(targetDims, strideOrder); } +// True when an axis of extent > 1 has stride 0, so several indices map to one element. +// The element span can still equal the element count when other axes leave matching gaps. +inline bool hasBroadcastAxis(const std::vector& dims, const std::vector& strides) +{ + for(size_t axis = 0; axis < dims.size(); ++axis) + { + if(dims[axis] > 1 && strides[axis] == 0) + { + return true; + } + } + return false; +} + // Checks if the tensor defined by dims and strides is packed (contiguous in memory). // Note: Assumes dims are positive (validated at graph level). // Strides can be negative (for reversed dimensions). @@ -306,14 +320,9 @@ inline bool isTensorPacked(const std::vector& dims, const std::vector 1 maps several indices to one element, and the - // span check below can still pass when other axes leave matching gaps. - for(size_t axis = 0; axis < dims.size(); ++axis) + if(hasBroadcastAxis(dims, strides)) { - if(dims[axis] > 1 && strides[axis] == 0) - { - return false; - } + return false; } // Calculate total element count diff --git a/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/Tensor.hpp b/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/Tensor.hpp index 0f2b1f19dd19..820fd0da1680 100644 --- a/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/Tensor.hpp +++ b/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/Tensor.hpp @@ -586,10 +586,10 @@ class ITensor virtual ITensorIterator cbegin() const = 0; virtual ITensorIterator cend() const = 0; - /// True iff the elements occupy the memory span with no gaps, i.e. - /// elementCount == elementSpace. This says nothing about the ORDER in which a - /// linear walk of that span visits them; a permutation of the strides is packed - /// too, and iterating one linearly visits the right addresses in the wrong order. + /// True iff elementCount == elementSpace and no dimension of size greater than 1 + /// has a zero stride. This says nothing about the ORDER in which a linear walk of + /// that span visits the elements; a permutation of the strides is packed too, and + /// iterating one linearly visits the right addresses in the wrong order. virtual bool isPacked() const = 0; virtual void markHostModified() = 0; @@ -837,14 +837,9 @@ class TensorBase : public ITensor bool computeIsPacked(const std::vector& dims, const std::vector& strides) const { - // A zero stride on an axis of extent > 1 maps several indices to one element. - // Its element space can still equal the item count, so check the strides first. - for(size_t axis = 0; axis < dims.size(); ++axis) + if(hasBroadcastAxis(dims, strides)) { - if(dims[axis] > 1 && strides[axis] == 0) - { - return false; - } + return false; } // Item count = largest stride * item count in that dimension return (calculateItemCount(dims) == calculateElementSpace(dims, strides)); @@ -886,8 +881,8 @@ class Tensor : public TensorBase , _elementCount(TensorBase::calculateItemCount(dims)) { validateDimsAndStridesSameSize(); - validateAllPositive(_dims, "dimension"); - validateAllNonNegative(_strides, "stride"); + validateAllAtLeast(_dims, 1, "positive", "dimension"); + validateAllAtLeast(_strides, 0, "non-negative", "stride"); // Set packed flag after validations since it can be incorrect if dims/strides are invalid. _packed = TensorBase::computeIsPacked(dims, strides); @@ -988,30 +983,18 @@ class Tensor : public TensorBase } } - void validateAllPositive(const std::vector& values, const std::string& valueName) const - { - for(size_t i = 0; i < values.size(); ++i) - { - if(values[i] <= 0) - { - std::ostringstream oss; - oss << "All " << valueName << "s must be positive. " << valueName << " " << i - << " is " << values[i]; - throw std::invalid_argument(oss.str()); - } - } - } - - void validateAllNonNegative(const std::vector& values, - const std::string& valueName) const + void validateAllAtLeast(const std::vector& values, + int64_t minimum, + const char* requirement, + const std::string& valueName) const { for(size_t i = 0; i < values.size(); ++i) { - if(values[i] < 0) + if(values[i] < minimum) { std::ostringstream oss; - oss << "All " << valueName << "s must be non-negative. " << valueName << " " << i - << " is " << values[i]; + oss << "All " << valueName << "s must be " << requirement << ". " << valueName + << " " << i << " is " << values[i]; throw std::invalid_argument(oss.str()); } } diff --git a/projects/hipdnn/data_sdk/tests/utilities/TestTensor.cpp b/projects/hipdnn/data_sdk/tests/utilities/TestTensor.cpp index 465f05a4e1b6..ddb716035c98 100644 --- a/projects/hipdnn/data_sdk/tests/utilities/TestTensor.cpp +++ b/projects/hipdnn/data_sdk/tests/utilities/TestTensor.cpp @@ -839,21 +839,30 @@ TEST(TestTensor, ZeroStrideOnUnitExtentAxisStaysPacked) TEST(TestTensor, FillWithRandomValuesBroadcastsAlongZeroStrideAxis) { + using hipdnn_data_sdk::types::isnan; + const std::vector dims = {3, 4, 2}; Tensor tensor(dims, {2, 0, 1}); + ASSERT_EQ(tensor.elementSpace(), 6); + tensor.fillWithSentinelValue(); tensor.fillWithRandomValues(-1.0F, 1.0F, 1234); + const auto* hostData = static_cast(tensor.rawHostData()); + for(size_t offset = 0; offset < tensor.elementSpace(); ++offset) + { + EXPECT_FALSE(isnan(hostData[offset])) << "offset " << offset; + EXPECT_GE(hostData[offset], -1.0F); + EXPECT_LE(hostData[offset], 1.0F); + } + for(int64_t i0 = 0; i0 < dims[0]; ++i0) { - for(int64_t i2 = 0; i2 < dims[2]; ++i2) + for(int64_t i1 = 0; i1 < dims[1]; ++i1) { - const float first = tensor.getHostValue(i0, 0, i2); - EXPECT_GE(first, -1.0F); - EXPECT_LE(first, 1.0F); - for(int64_t i1 = 1; i1 < dims[1]; ++i1) + for(int64_t i2 = 0; i2 < dims[2]; ++i2) { - EXPECT_EQ(tensor.getHostValue(i0, i1, i2), first); + EXPECT_EQ(tensor.getHostValue(i0, i1, i2), hostData[tensor.getIndex(i0, i1, i2)]); } } } From 86bb42199247dfb5ce005ef0aae5ff7b6afd9052 Mon Sep 17 00:00:00 2001 From: CMiservaAMD Date: Fri, 30 Oct 2026 18:56:43 -0400 Subject: [PATCH 3/3] (CMCL) hipDNN data SDK: drop redundant zero-stride test assertions and update isTensorPacked comment Remove IsTensorPackedZeroStrideBroadcast, whose cases were already rejected by the span check, and the identity loop in the zero-stride fill test. State the zero-stride rule in the isTensorPacked comment instead of the claim that negative strides are supported. ALMIOPEN-3109 Co-Authored-By: Claude Opus 5.5 (1M context) --- .../hipdnn_data_sdk/utilities/ShapeUtilities.hpp | 3 +-- .../data_sdk/tests/utilities/TestShapeUtilities.cpp | 6 ------ .../hipdnn/data_sdk/tests/utilities/TestTensor.cpp | 11 ----------- 3 files changed, 1 insertion(+), 19 deletions(-) diff --git a/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/ShapeUtilities.hpp b/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/ShapeUtilities.hpp index 3fb168cd50a7..214d382259af 100644 --- a/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/ShapeUtilities.hpp +++ b/projects/hipdnn/data_sdk/include/hipdnn_data_sdk/utilities/ShapeUtilities.hpp @@ -305,8 +305,7 @@ inline bool hasBroadcastAxis(const std::vector& dims, const std::vector } // Checks if the tensor defined by dims and strides is packed (contiguous in memory). -// Note: Assumes dims are positive (validated at graph level). -// Strides can be negative (for reversed dimensions). +// Assumes dims are positive. A zero stride on a dimension of size > 1 (broadcast) is never packed. inline bool isTensorPacked(const std::vector& dims, const std::vector& strides) { if(dims.size() != strides.size()) diff --git a/projects/hipdnn/data_sdk/tests/utilities/TestShapeUtilities.cpp b/projects/hipdnn/data_sdk/tests/utilities/TestShapeUtilities.cpp index b3711c0bd304..5d73b86e0460 100644 --- a/projects/hipdnn/data_sdk/tests/utilities/TestShapeUtilities.cpp +++ b/projects/hipdnn/data_sdk/tests/utilities/TestShapeUtilities.cpp @@ -783,12 +783,6 @@ TEST(TestShapeUtils, IsTensorPackedZeroStrideWithMatchingSpan) EXPECT_FALSE(isTensorPacked({2, 2}, {3, 0})); } -TEST(TestShapeUtils, IsTensorPackedZeroStrideBroadcast) -{ - EXPECT_FALSE(isTensorPacked({2, 3, 4}, {4, 0, 1})); - EXPECT_FALSE(isTensorPacked({2, 3}, {0, 0})); -} - TEST(TestShapeUtils, IsTensorPackedZeroStrideOnUnitExtentAxis) { EXPECT_TRUE(isTensorPacked({2, 1, 3}, {3, 0, 1})); diff --git a/projects/hipdnn/data_sdk/tests/utilities/TestTensor.cpp b/projects/hipdnn/data_sdk/tests/utilities/TestTensor.cpp index ddb716035c98..94a6bf2f2ff2 100644 --- a/projects/hipdnn/data_sdk/tests/utilities/TestTensor.cpp +++ b/projects/hipdnn/data_sdk/tests/utilities/TestTensor.cpp @@ -855,17 +855,6 @@ TEST(TestTensor, FillWithRandomValuesBroadcastsAlongZeroStrideAxis) EXPECT_GE(hostData[offset], -1.0F); EXPECT_LE(hostData[offset], 1.0F); } - - for(int64_t i0 = 0; i0 < dims[0]; ++i0) - { - for(int64_t i1 = 0; i1 < dims[1]; ++i1) - { - for(int64_t i2 = 0; i2 < dims[2]; ++i2) - { - EXPECT_EQ(tensor.getHostValue(i0, i1, i2), hostData[tensor.getIndex(i0, i1, i2)]); - } - } - } } TEST(TestTensor, ZeroStrideTensorRoundTripsThroughDevice)