From 7effb61bb81def3c16d95b7ccdca2e4e6bc635e6 Mon Sep 17 00:00:00 2001 From: Zach Vincze Date: Fri, 3 Jul 2026 15:02:53 -0400 Subject: [PATCH 1/5] Zero-initialize shapes and strides arrays --- include/core/tensor_shape.hpp | 2 +- python/src/py_tensor.cpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/include/core/tensor_shape.hpp b/include/core/tensor_shape.hpp index 68e8f8d8..c6480026 100644 --- a/include/core/tensor_shape.hpp +++ b/include/core/tensor_shape.hpp @@ -124,7 +124,7 @@ class TensorShape { bool operator!=(const TensorShape &rhs) const; private: - std::array m_shape; + std::array m_shape{}; TensorLayout m_layout; size_t m_size; }; diff --git a/python/src/py_tensor.cpp b/python/src/py_tensor.cpp index 03cebdf4..23e143d8 100644 --- a/python/src/py_tensor.cpp +++ b/python/src/py_tensor.cpp @@ -108,7 +108,7 @@ std::shared_ptr PyTensor::fromDLPack(pybind11::object src, eTensorLayo dlpackCapsule.set_name("used_dltensor"); // Copy the shape data from DLPack to a fixed-size array - std::array shapeData; + std::array shapeData{}; for (int i = 0; i < dlTensor.ndim; ++i) { shapeData[i] = dlTensor.shape[i]; } @@ -117,7 +117,7 @@ std::shared_ptr PyTensor::fromDLPack(pybind11::object src, eTensorLayo eDataType dtype = DLTypeToRoccvType(dlTensor.dtype); // Prepare the strides array - std::array stridesData; + std::array stridesData{}; // If strides are not present, assume contiguous layout. We really shouldn't be recalculating strides here. DLPack // now enforces that the strides are present, but we'll keep this for backwards compatibility. From d8d377517eb4310df8d25803f6b71514bc298e47 Mon Sep 17 00:00:00 2001 From: Zach Vincze Date: Fri, 3 Jul 2026 15:16:02 -0400 Subject: [PATCH 2/5] Revert zero-initialization and use already implemented TensorShape == operator --- include/core/tensor_shape.hpp | 2 +- python/src/py_tensor.cpp | 4 ++-- src/core/tensor.cpp | 3 +-- 3 files changed, 4 insertions(+), 5 deletions(-) diff --git a/include/core/tensor_shape.hpp b/include/core/tensor_shape.hpp index c6480026..68e8f8d8 100644 --- a/include/core/tensor_shape.hpp +++ b/include/core/tensor_shape.hpp @@ -124,7 +124,7 @@ class TensorShape { bool operator!=(const TensorShape &rhs) const; private: - std::array m_shape{}; + std::array m_shape; TensorLayout m_layout; size_t m_size; }; diff --git a/python/src/py_tensor.cpp b/python/src/py_tensor.cpp index 23e143d8..03cebdf4 100644 --- a/python/src/py_tensor.cpp +++ b/python/src/py_tensor.cpp @@ -108,7 +108,7 @@ std::shared_ptr PyTensor::fromDLPack(pybind11::object src, eTensorLayo dlpackCapsule.set_name("used_dltensor"); // Copy the shape data from DLPack to a fixed-size array - std::array shapeData{}; + std::array shapeData; for (int i = 0; i < dlTensor.ndim; ++i) { shapeData[i] = dlTensor.shape[i]; } @@ -117,7 +117,7 @@ std::shared_ptr PyTensor::fromDLPack(pybind11::object src, eTensorLayo eDataType dtype = DLTypeToRoccvType(dlTensor.dtype); // Prepare the strides array - std::array stridesData{}; + std::array stridesData; // If strides are not present, assume contiguous layout. We really shouldn't be recalculating strides here. DLPack // now enforces that the strides are present, but we'll keep this for backwards compatibility. diff --git a/src/core/tensor.cpp b/src/core/tensor.cpp index b8ae2c2f..3950b7cd 100644 --- a/src/core/tensor.cpp +++ b/src/core/tensor.cpp @@ -387,8 +387,7 @@ void Tensor::copyToHostAsync(void* dst, hipStream_t stream) const { void Tensor::copyToAsync(const Tensor& dst, hipStream_t stream) const { // Source and destination must describe the same logical data; only their padding (row pitch) may differ. - if (rank() != dst.rank() || m_requirements.shape != dst.m_requirements.shape || - dtype().size() != dst.dtype().size()) { + if (shape() != dst.shape() || dtype().size() != dst.dtype().size()) { throw Exception("Source and destination tensors must have the same shape and element size for a copy.", eStatusType::INVALID_VALUE); } From 50ea60db4973ac77758c2c37038d0c1ab096304e Mon Sep 17 00:00:00 2001 From: Zach Vincze Date: Fri, 3 Jul 2026 15:22:01 -0400 Subject: [PATCH 3/5] Add regression tests for tensor-to-tensor copies --- .../cpp/src/tests/core/tensor/test_tensor.cpp | 68 +++++++++++++++++++ 1 file changed, 68 insertions(+) diff --git a/tests/roccv/cpp/src/tests/core/tensor/test_tensor.cpp b/tests/roccv/cpp/src/tests/core/tensor/test_tensor.cpp index 2f3ec4b3..14020688 100644 --- a/tests/roccv/cpp/src/tests/core/tensor/test_tensor.cpp +++ b/tests/roccv/cpp/src/tests/core/tensor/test_tensor.cpp @@ -197,6 +197,65 @@ void TestTensorCopyCorrectness() { EXPECT_VECTOR_EQ(inputDataHost, outputDataHost); } +/** + * @brief Tests correctness of tensor-to-tensor copies (Tensor::copyToAsync). + * + * Regression coverage for the shape-equality check in copyToAsync: source and destination describe the same logical + * shape but may carry different row padding (e.g. a padded GPU tensor copied into a packed CPU tensor). The copy must + * succeed and preserve data. This guards against comparing unused trailing shape dimensions, which previously caused + * otherwise-identical tensors to be rejected. + * + * @param[in] srcShape The logical shape shared by source and destination. + * @param[in] dtype The datatype shared by source and destination. + * @param[in] srcDevice Device the source tensor is allocated on. + * @param[in] dstDevice Device the destination tensor is allocated on. + */ +void TestTensorCopyToCorrectness(const TensorShape& srcShape, const DataType& dtype, eDeviceType srcDevice, + eDeviceType dstDevice) { + Tensor source(srcShape, dtype, srcDevice); + Tensor dest(srcShape, dtype, dstDevice); + + const size_t hostDataSize = source.shape().size() * dtype.size(); + std::vector inputDataHost(hostDataSize); + for (size_t i = 0; i < inputDataHost.size(); i++) { + inputDataHost[i] = static_cast(i % 256); + } + + hipStream_t stream; + HIP_VALIDATE_NO_ERRORS(hipStreamCreate(&stream)); + + // Seed the source, copy source -> dest, then read the destination back to the host for comparison. + source.copyFromHostAsync(inputDataHost.data(), stream); + source.copyToAsync(dest, stream); + + std::vector outputDataHost(hostDataSize); + dest.copyToHostAsync(outputDataHost.data(), stream); + + HIP_VALIDATE_NO_ERRORS(hipStreamSynchronize(stream)); + HIP_VALIDATE_NO_ERRORS(hipStreamDestroy(stream)); + + EXPECT_VECTOR_EQ(inputDataHost, outputDataHost); +} + +/** + * @brief Negative tests for Tensor::copyToAsync, verifying mismatched source/destination are rejected. + */ +void TestNegativeTensorCopyTo() { + Tensor source(TensorShape({2, 4, 4}, "HWC"), DataType(DATA_TYPE_U8)); + + // Differing dimensions must be rejected. + { + Tensor dest(TensorShape({2, 4, 3}, "HWC"), DataType(DATA_TYPE_U8)); + EXPECT_EXCEPTION(source.copyToAsync(dest), eStatusType::INVALID_VALUE); + } + + // Differing element size must be rejected. + { + Tensor dest(TensorShape({2, 4, 4}, "HWC"), DataType(DATA_TYPE_U16)); + EXPECT_EXCEPTION(source.copyToAsync(dest), eStatusType::INVALID_VALUE); + } +} + /** * @brief Tests internal stride calculations on Tensor construction. */ @@ -231,12 +290,21 @@ int main(int argc, char** argv) { TEST_CASE(TestNegativeTensorShape()); TEST_CASE(TestNegativeTensor()); TEST_CASE(TestNegativeTensorReshape()); + TEST_CASE(TestNegativeTensorCopyTo()); // Correctness tests TEST_CASE(TestTensorCorrectness()); TEST_CASE(TestTensorReshapeCorrectness()); TEST_CASE(TestTensorCopyCorrectness()); + // Tensor-to-tensor copy tests (padded/cross-device round-trips) + TEST_CASE(TestTensorCopyToCorrectness(TensorShape({2, 10, 10, 3}, "NHWC"), DataType(DATA_TYPE_U8), eDeviceType::GPU, + eDeviceType::CPU)); + TEST_CASE(TestTensorCopyToCorrectness(TensorShape({2, 10, 10, 3}, "NHWC"), DataType(DATA_TYPE_U8), eDeviceType::GPU, + eDeviceType::GPU)); + TEST_CASE(TestTensorCopyToCorrectness(TensorShape({4, 100}, "NW"), DataType(DATA_TYPE_F32), eDeviceType::CPU, + eDeviceType::GPU)); + // Stride calculation tests // clang-format off TEST_CASE(TestTensorStrideCalculation(TensorShape({1, 2, 4, 4}, "NHWC"), DataType(DATA_TYPE_U8))); From bed580eb464f78ae766751725014016bf5a2d175 Mon Sep 17 00:00:00 2001 From: Zach Vincze Date: Fri, 3 Jul 2026 15:27:19 -0400 Subject: [PATCH 4/5] Remove redundant comment --- tests/roccv/cpp/src/tests/core/tensor/test_tensor.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/tests/roccv/cpp/src/tests/core/tensor/test_tensor.cpp b/tests/roccv/cpp/src/tests/core/tensor/test_tensor.cpp index 14020688..3c833c54 100644 --- a/tests/roccv/cpp/src/tests/core/tensor/test_tensor.cpp +++ b/tests/roccv/cpp/src/tests/core/tensor/test_tensor.cpp @@ -202,8 +202,7 @@ void TestTensorCopyCorrectness() { * * Regression coverage for the shape-equality check in copyToAsync: source and destination describe the same logical * shape but may carry different row padding (e.g. a padded GPU tensor copied into a packed CPU tensor). The copy must - * succeed and preserve data. This guards against comparing unused trailing shape dimensions, which previously caused - * otherwise-identical tensors to be rejected. + * succeed and preserve data. * * @param[in] srcShape The logical shape shared by source and destination. * @param[in] dtype The datatype shared by source and destination. From 35c34c846743d7d7efa29b831fde4f8f4f9252b0 Mon Sep 17 00:00:00 2001 From: Zach Vincze Date: Fri, 3 Jul 2026 15:37:56 -0400 Subject: [PATCH 5/5] Poison trailing shape data to effectively guard against regression for copy to --- .../roccv/cpp/src/tests/core/tensor/test_tensor.cpp | 13 +++++++++++-- 1 file changed, 11 insertions(+), 2 deletions(-) diff --git a/tests/roccv/cpp/src/tests/core/tensor/test_tensor.cpp b/tests/roccv/cpp/src/tests/core/tensor/test_tensor.cpp index 3c833c54..32be8be1 100644 --- a/tests/roccv/cpp/src/tests/core/tensor/test_tensor.cpp +++ b/tests/roccv/cpp/src/tests/core/tensor/test_tensor.cpp @@ -211,8 +211,17 @@ void TestTensorCopyCorrectness() { */ void TestTensorCopyToCorrectness(const TensorShape& srcShape, const DataType& dtype, eDeviceType srcDevice, eDeviceType dstDevice) { - Tensor source(srcShape, dtype, srcDevice); - Tensor dest(srcShape, dtype, dstDevice); + // Build requirements normally, then poison the unused trailing shape slots with distinct values per tensor. The + // allocation size and strides derive from the in-rank dimensions, so this does not affect the actual data layout. + Tensor::Requirements srcReqs = Tensor::CalcRequirements(srcShape, dtype, srcDevice); + Tensor::Requirements dstReqs = Tensor::CalcRequirements(srcShape, dtype, dstDevice); + for (int i = srcReqs.rank; i < ROCCV_TENSOR_MAX_RANK; i++) { + srcReqs.shape[i] = 100 + i; + dstReqs.shape[i] = 900 - i; + } + + Tensor source(srcReqs); + Tensor dest(dstReqs); const size_t hostDataSize = source.shape().size() * dtype.size(); std::vector inputDataHost(hostDataSize);