From 0640152cfe8f936ad607e0060dd0fa786e4adbcc Mon Sep 17 00:00:00 2001 From: "Ding, Yi" Date: Wed, 7 Oct 2026 23:49:53 -0400 Subject: [PATCH] fix(tensilelite): copy the source tensor in CopyTensorVoid CopyTensorVoid took its source address from dst, so every hipMemcpyAsync copied the destination onto itself. Its only caller is copyBadInputBuffers, which fills a NaN bounds-check buffer with sentinels and then relies on CopyTensorVoid to write the real tensor into the middle. The tensor never landed, so with --bounds-check=NaN every input and output kept the inf sentinel: - the kernel and the CPU reference both computed D from inf inputs, so validation compared inf against inf and passed for any result; - the Synchronizer buffer started as inf instead of zero, so the Synchronizer clean-on-exit check reported StreamK solutions as dirty. Measured with the tensilelite client on gfx950, BF16 GEMM, 3 sizes x 8 solutions, --bounds-check=NaN --print-valids: - before: every compared element printed inf==inf; 12/24 FAILED, all from the Synchronizer check - after: compared elements print ordinary values; 24/24 PASSED - Disable mode: 24/24 PASSED before and after --- .../include/Tensile/hip/HipUtils.hpp | 2 +- .../tensilelite/tests/DataInit_test.cpp | 18 ++++++++++++++++++ 2 files changed, 19 insertions(+), 1 deletion(-) diff --git a/projects/hipblaslt/tensilelite/include/Tensile/hip/HipUtils.hpp b/projects/hipblaslt/tensilelite/include/Tensile/hip/HipUtils.hpp index ad12487f6b19..1fb8bcda962c 100644 --- a/projects/hipblaslt/tensilelite/include/Tensile/hip/HipUtils.hpp +++ b/projects/hipblaslt/tensilelite/include/Tensile/hip/HipUtils.hpp @@ -198,7 +198,7 @@ namespace TensileLite auto beginOffset = desc.index(coord); size_t bytesOffset = multiplyElementSize(beginOffset, desc.elementBytes()); uint8_t* dstBytes = (uint8_t*)dst + bytesOffset; - uint8_t* srcBytes = (uint8_t*)dst + bytesOffset; + uint8_t* srcBytes = (uint8_t*)src + bytesOffset; HIP_CHECK_EXC(hipMemcpyAsync(dstBytes, srcBytes, copyBytes, direction, stream)); } diff --git a/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp b/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp index 53f00f5dc121..268ca76a7ce3 100644 --- a/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp +++ b/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp @@ -13,6 +13,7 @@ #include #include #include +#include #include #include #include @@ -20,6 +21,7 @@ #include #include #include +#include #include #include #include @@ -563,3 +565,19 @@ TEST(DecodeMXElement, UnsupportedDataTypeReturnsNaN) } #endif // HIPBLASLT_ENABLE_MXDATAGENERATOR + +// ============================================================================= +// Section 5 — TensileLite::hip::CopyTensorVoid +// ============================================================================= +TEST(CopyTensorVoid, CopiesSource) +{ + TensorDescriptor desc("t", rocisa::DataType::Float, {3, 2}); + std::vector src(desc.totalAllocatedElements()); + std::iota(src.begin(), src.end(), 1.0f); + std::vector dst(src.size(), 0.0f); + + TensileLite::hip::CopyTensorVoid(dst.data(), src.data(), desc, hipMemcpyHostToHost); + ASSERT_EQ(hipStreamSynchronize(nullptr), hipSuccess); + + EXPECT_EQ(dst, src); +}