CMiservaAMD · Nov 2, 2026 · Oct 29, 2026 · Oct 29, 2026 · Oct 30, 2026 · Nov 2, 2026
@@ -290,6 +290,20 @@ inline std::vector<int64_t>
     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<int64_t>& dims, const std::vector<int64_t>& 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,6 +320,11 @@ inline bool isTensorPacked(const std::vector<int64_t>& dims, const std::vector<i
         return true;
     }
 
+    if(hasBroadcastAxis(dims, strides))
+    {
+        return false;
+    }
+
     // Calculate total element count
     const auto count
         = std::accumulate(dims.begin(), dims.end(), static_cast<int64_t>(1), std::multiplies<>());

@@ -586,10 +586,10 @@ class ITensor
     virtual ITensorIterator<true> cbegin() const = 0;
     virtual ITensorIterator<true> 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,6 +837,10 @@ class TensorBase : public ITensor
     bool computeIsPacked(const std::vector<int64_t>& dims,
                          const std::vector<int64_t>& strides) const
     {
+        if(hasBroadcastAxis(dims, strides))
+        {
+            return false;
+        }
         // Item count = largest stride * item count in that dimension
         return (calculateItemCount(dims) == calculateElementSpace(dims, strides));
     }
@@ -877,8 +881,8 @@ class Tensor : public TensorBase<T>
         , _elementCount(TensorBase<T>::calculateItemCount(dims))
     {
         validateDimsAndStridesSameSize();
-        validateAllPositive(_dims, "dimension");
-        validateAllPositive(_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<T>::computeIsPacked(dims, strides);
@@ -979,15 +983,18 @@ class Tensor : public TensorBase<T>
         }
     }
 
-    void validateAllPositive(const std::vector<int64_t>& values, const std::string& valueName) const
+    void validateAllAtLeast(const std::vector<int64_t>& 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 positive. " << valueName << " " << i
-                    << " is " << values[i];
+                oss << "All " << valueName << "s must be " << requirement << ". " << valueName
+                    << " " << i << " is " << values[i];
                 throw std::invalid_argument(oss.str());
             }
         }

@@ -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

@@ -790,12 +790,99 @@ TEST(TestTensor, ConstructorThrowsOnNegativeStride)
     EXPECT_THROW(const Tensor<float> tensor(dims, strides), std::invalid_argument);
 }
 
-TEST(TestTensor, ConstructorThrowsOnZeroStride)
+TEST(TestTensor, ZeroStrideAxisReadsOneElement)
 {
     const std::vector<int64_t> dims = {2, 3, 4};
-    const std::vector<int64_t> strides = {12, 0, 1};
+    const std::vector<int64_t> strides = {4, 0, 1};
 
-    EXPECT_THROW(const Tensor<float> tensor(dims, strides), std::invalid_argument);
+    Tensor<float> 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<float> 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<float> tensor({2, 2}, {3, 0});
+
+    ASSERT_EQ(tensor.elementCount(), tensor.elementSpace());
+    EXPECT_FALSE(tensor.isPacked());
+}
+
+TEST(TestTensor, ZeroStrideOnUnitExtentAxisStaysPacked)
+{
+    const Tensor<float> tensor({2, 1, 3}, {3, 0, 1});
+
+    EXPECT_TRUE(tensor.isPacked());
+}
+
+TEST(TestTensor, FillWithRandomValuesBroadcastsAlongZeroStrideAxis)
+{
+    using hipdnn_data_sdk::types::isnan;
+
+    const std::vector<int64_t> dims = {3, 4, 2};
+    Tensor<float> tensor(dims, {2, 0, 1});
+    ASSERT_EQ(tensor.elementSpace(), 6);
+
+    tensor.fillWithSentinelValue();
+    tensor.fillWithRandomValues(-1.0F, 1.0F, 1234);
+
+    const auto* hostData = static_cast<const float*>(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 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)
+{
+    SKIP_IF_NO_DEVICES();
+
+    Tensor<float> tensor({2, 3}, {1, 0});
+    tensor.setHostValue(1.5F, 0, 0);
+    tensor.setHostValue(2.5F, 1, 0);
+
+    auto* device = tensor.rawDeviceData();
+    std::vector<float> fromDevice(tensor.elementSpace());
+    ASSERT_EQ(
+        hipMemcpy(
+            fromDevice.data(), device, fromDevice.size() * sizeof(float), hipMemcpyDeviceToHost),
+        hipSuccess);
+    EXPECT_EQ(fromDevice, (std::vector<float>{1.5F, 2.5F}));
 }
 
 TEST(TestTensor, ConstructorThrowsOnMultipleInvalidDimensions)