From 48dd3e5f7780f05b55494d7c7de4b2704906a176 Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Fri, 2 Oct 2026 13:01:49 -0400 Subject: [PATCH 1/8] test(hipblaslt): AIHPBLAS-4994 place fast_check operands across 4 GiB boundaries JIRA ID : AIHPBLAS-4994 Some address defects depend on where a buffer sits rather than on the problem size. ROCM-32046 drops the carry out of the low 32 bits of an address, so it goes wrong only when an operand crosses a 4 GiB boundary, and the wrong access lands 4 GiB away, outside every allocation. fast_check's padding poison and D sentinel cannot see such an access. On gfx1250 the page fault it causes can leave the GPU unusable (ROCM-32049). - A new placement test argument names one operand (a, b, c, d, bias, scale_alpha_vec or workspace), plus an offset. PlacedRegion, in hip_placement.hpp, uses HIP virtual memory management to map that operand across a 4 GiB boundary, with poison windows 4 GiB below and above it. A dropped-carry read returns the poison value, so fast_check reports a wrong D instead of the GPU faulting. A dropped-carry write changes the poison, which is checked after every solution. - HipDeviceBuffer gains a placement constructor. Its move operations now clear the moved-from pointer, so a moved-from buffer no longer re-checks guard bytes in memory it no longer owns. - Self-tests use a kernel that drops the carry on purpose, plus a placed buffer moved by vector growth. - New YAML cases place each operand in turn on a small shape. The pre_checkin cases run the heuristic's first 16 solutions and the nightly cases run every solution. Co-Authored-By: Claude Opus 5.5 --- projects/hipblaslt/TESTING.md | 2 +- .../clients/common/include/fast_check.hpp | 18 ++ .../clients/common/include/hipBuffer.hpp | 59 ++++- .../clients/common/include/hip_placement.hpp | 212 ++++++++++++++++++ .../common/include/hipblaslt_arguments.hpp | 6 + .../clients/common/include/testing_matmul.hpp | 159 +++++++++++-- .../clients/common/src/fast_check.cpp | 119 +++++++--- .../common/src/hipblaslt_arguments.cpp | 2 + .../clients/tests/data/hipblaslt_common.yaml | 4 + .../clients/tests/data/matmul_gtest.yaml | 105 +++++++++ .../clients/tests/src/fast_check_gtest.cpp | 121 ++++++++++ .../clients/tests/src/matmul_gtest.cpp | 2 + 12 files changed, 755 insertions(+), 54 deletions(-) create mode 100644 projects/hipblaslt/clients/common/include/hip_placement.hpp diff --git a/projects/hipblaslt/TESTING.md b/projects/hipblaslt/TESTING.md index 28e441420155..c2310f5c5940 100644 --- a/projects/hipblaslt/TESTING.md +++ b/projects/hipblaslt/TESTING.md @@ -244,7 +244,7 @@ There is no separate unit-test binary. The host-only tests that do exist are com | `caching_library_gtest.cpp` | `CachingLibraryCollision` | Solution-cache keying under a forced hash collision | | `ulp_gtest.cpp` | `UlpMantissaBits`, `UlpAsDouble`, `UlpDistance`, `UlpAccumulate`, `UlpCheckGeneral` | Units-in-the-last-place error-measurement helpers, pure math with no GPU dependency | | `unit_gtest.cpp` | `UnitCheckIdentical` | Exact-storage checks used to skip element-by-element output comparisons when the addressed values have identical bytes | -| `fast_check_gtest.cpp` | `FastCheck_pre_checkin`, `FastCheckDevice_pre_checkin` | The `fast_check` verifier: fault injection into host-computed results (`FastCheck`, no GPU), and the device padding, poison and probe-sum kernels on small buffers (`FastCheckDevice`, no hipBLASLt kernels) | +| `fast_check_gtest.cpp` | `FastCheck_pre_checkin`, `FastCheckDevice_pre_checkin` | The `fast_check` verifier: fault injection into host-computed results (`FastCheck_pre_checkin`, no GPU), and the device padding, poison and probe-sum kernels plus 4 GiB-boundary buffer placement on small buffers (`FastCheckDevice_pre_checkin`, no hipBLASLt kernels) | | `arch_revision_gtest.cpp` | `ArchRevisionSmoke` | ASIC-revision-to-library-arch mapping, a pure function reached by relative `#include` of an internal header (see point 1 below) | That is a handful of test cases of genuinely hardware-independent logic, against a library of diff --git a/projects/hipblaslt/clients/common/include/fast_check.hpp b/projects/hipblaslt/clients/common/include/fast_check.hpp index d991b7f9836a..69b1e7c37061 100644 --- a/projects/hipblaslt/clients/common/include/fast_check.hpp +++ b/projects/hipblaslt/clients/common/include/fast_check.hpp @@ -153,6 +153,24 @@ std::string fast_check_describe_buffers(const std::vector& buff // memory; total_elements is the size of its allocation in elements, which may extend past the // last batch. +// Bytes per element, or 0 when fast_check does not support the type. +size_t fast_check_element_size(hipDataType type); + +// The bit pattern of kFastCheckPoisonValue in the given type, in the low bytes. +uint64_t fast_check_poison_bits(hipDataType type); + +// Counts the elements of a device range that no longer hold kFastCheckPoisonValue. +struct FastCheckChanged +{ + uint64_t count = 0; + uint64_t first = 0; // index of the first changed element, when count > 0 + bool ok = true; // false when the scan itself failed +}; +FastCheckChanged fast_check_count_changed_device(const void* data, + hipDataType type, + size_t elements, + hipStream_t stream); + // Fills every element of a device buffer with fast_check_sentinel_bits(type). Returns the launch // error, or hipErrorInvalidValue for a type fast_check does not support. hipError_t fast_check_fill_sentinel_device(void* buffer, diff --git a/projects/hipblaslt/clients/common/include/hipBuffer.hpp b/projects/hipblaslt/clients/common/include/hipBuffer.hpp index 317a268e2ca7..ea07ca212901 100644 --- a/projects/hipblaslt/clients/common/include/hipBuffer.hpp +++ b/projects/hipblaslt/clients/common/include/hipBuffer.hpp @@ -28,8 +28,13 @@ #include "d_vector.hpp" #include "datatype_interface.hpp" +#include "hip_placement.hpp" #include "hipblaslt_ostream.hpp" +#include +#include +#include + class HipDeviceBuffer : public d_vector_type { public: @@ -40,16 +45,59 @@ class HipDeviceBuffer : public d_vector_type { } + // Places the buffer across a 4 GiB virtual-address boundary, offset_below bytes after its + // start, with poison memory mapped 4 GiB below and above it (see hip_placement.hpp). buf() + // is null when placement fails; *why and *unsupported then say why. + HipDeviceBuffer(hipDataType dtype, + std::size_t numElements, + std::size_t offset_below, + std::string* why, + bool* unsupported) + : d_vector_type(dtype, numElements, false) + , numBytes(realDataTypeSize(dtype) * numElements) + , placed(PlacedRegion::create(numBytes, offset_below, dtype, why, unsupported)) + , buffer(placed ? placed->ptr() : nullptr) + { + } + ~HipDeviceBuffer() { - this->device_vector_teardown(static_cast(buffer)); + // A placed buffer has no guard pads, and its memory belongs to `placed`. + this->device_vector_teardown(placed ? nullptr : static_cast(buffer)); buffer = nullptr; } HipDeviceBuffer(const HipDeviceBuffer&) = delete; - HipDeviceBuffer(HipDeviceBuffer&&) = default; HipDeviceBuffer& operator=(const HipDeviceBuffer&) = delete; - HipDeviceBuffer& operator=(HipDeviceBuffer&&) = default; + + // The moved-from buffer must forget its pointer, or its destructor would check guard pads + // in memory it no longer owns, which for a placed buffer is not mapped. + HipDeviceBuffer(HipDeviceBuffer&& other) noexcept + : d_vector_type(std::move(other)) + , numBytes(other.numBytes) + , placed(std::move(other.placed)) + , buffer(std::exchange(other.buffer, nullptr)) + { + } + + HipDeviceBuffer& operator=(HipDeviceBuffer&& other) noexcept + { + if(this != &other) + { + this->device_vector_teardown(placed ? nullptr : static_cast(buffer)); + d_vector_type::operator=(std::move(other)); + numBytes = other.numBytes; + placed = std::move(other.placed); + buffer = std::exchange(other.buffer, nullptr); + } + return *this; + } + + // The placement of this buffer, or null when it was allocated normally. + const PlacedRegion* placement() const + { + return placed.get(); + } void* buf() { @@ -84,8 +132,9 @@ class HipDeviceBuffer : public d_vector_type } private: - std::size_t numBytes; - void* buffer; + std::size_t numBytes; + std::unique_ptr placed; + void* buffer; }; class HipHostBuffer diff --git a/projects/hipblaslt/clients/common/include/hip_placement.hpp b/projects/hipblaslt/clients/common/include/hip_placement.hpp new file mode 100644 index 000000000000..eda6fccaa4d9 --- /dev/null +++ b/projects/hipblaslt/clients/common/include/hip_placement.hpp @@ -0,0 +1,212 @@ +// Copyright Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +// Placement of a device buffer across a 4 GiB virtual-address boundary, for testing address +// arithmetic that drops the carry out of the low 32 bits of an address (ROCM-32046). Such a +// defect goes wrong only when a buffer crosses a 4 GiB boundary, and the wrong access lands +// exactly 4 GiB away, outside every allocation. +// +// PlacedRegion reserves address space with HIP virtual memory management and maps three windows +// of the same size: the buffer, starting a chosen number of bytes before a 4 GiB boundary, and +// two poison windows exactly 4 GiB below and above it, filled with kFastCheckPoisonValue. A read +// whose carry is dropped lands in the low window and returns the poison value instead of +// faulting, so the GEMM result is wrong rather than the GPU stopping. A write whose carry is +// dropped changes the poison, which verify_poison() reports. + +#pragma once + +#include "fast_check.hpp" + +#include + +#include +#include +#include +#include + +class PlacedRegion +{ +public: + static constexpr uint64_t k4GiB = uint64_t(1) << 32; + + // Maps `bytes` so that the 4 GiB boundary falls `offset_below` bytes after the start. + // offset_below = 0 picks the middle of the buffer, rounded down to the mapping granularity. + // Returns nullptr and sets *why on failure; *unsupported is true when the GPU or platform + // cannot do this at all, as opposed to an invalid request. + static std::unique_ptr create( + size_t bytes, size_t offset_below, hipDataType type, std::string* why, bool* unsupported) + { + auto fail = [&](const std::string& text, bool platform) { + if(why) + *why = text; + if(unsupported) + *unsupported = platform; + return std::unique_ptr(); + }; +#ifdef _WIN32 + return fail("buffer placement needs HIP virtual memory management, which is not " + "supported on Windows", + true); +#else + int device = 0, vmm = 0; + if(hipGetDevice(&device) != hipSuccess + || hipDeviceGetAttribute( + &vmm, hipDeviceAttributeVirtualMemoryManagementSupported, device) + != hipSuccess + || !vmm) + return fail("this GPU does not support HIP virtual memory management", true); + + const size_t es = fast_check_element_size(type); + if(!es) + return fail("buffer placement does not support data type " + std::to_string(int(type)), + false); + + std::unique_ptr r(new PlacedRegion); + r->m_type = type; + r->m_prop.type = hipMemAllocationTypePinned; + r->m_prop.location.type = hipMemLocationTypeDevice; + r->m_prop.location.id = device; + size_t granularity = 0; + if(hipMemGetAllocationGranularity( + &granularity, &r->m_prop, hipMemAllocationGranularityMinimum) + != hipSuccess + || !granularity || granularity % es) + return fail("could not get a usable mapping granularity", true); + + if(offset_below == 0) + offset_below = std::max(granularity, bytes / 2 / granularity * granularity); + if(offset_below % granularity) + return fail("the placement offset must be a multiple of the mapping granularity (" + + std::to_string(granularity) + " bytes)", + false); + if(offset_below >= bytes) + return fail("a " + std::to_string(bytes) + "-byte buffer cannot cross a 4 GiB boundary " + + std::to_string(offset_below) + " bytes after its start", + false); + + r->m_bytes = bytes; + r->m_span = (bytes + granularity - 1) / granularity * granularity; + if(r->m_span >= k4GiB) + return fail("a placed buffer must be smaller than 4 GiB", false); + + // Reserve enough to put both poison windows around a boundary inside the reservation, + // and compute the boundary from the address actually returned; the address hint is not + // guaranteed. + r->m_reserved_bytes = 4 * k4GiB + 2 * r->m_span; + if(hipMemAddressReserve(&r->m_reserved, r->m_reserved_bytes, 0, nullptr, 0) != hipSuccess) + { + r->m_reserved = nullptr; + return fail("could not reserve address space for placement", false); + } + const uint64_t base = reinterpret_cast(r->m_reserved); + const uint64_t boundary = (base + k4GiB + r->m_span + k4GiB - 1) / k4GiB * k4GiB; + r->m_windows[0] = boundary - offset_below; // the buffer + r->m_windows[1] = r->m_windows[0] - k4GiB; // poison below + r->m_windows[2] = r->m_windows[0] + k4GiB; // poison above + + hipMemAccessDesc access{}; + access.location = r->m_prop.location; + access.flags = hipMemAccessFlagsProtReadWrite; + for(int w = 0; w < 3; w++) + { + // Each window has its own handle: a mapping must cover a whole handle. + void* at = reinterpret_cast(r->m_windows[w]); + if(hipMemCreate(&r->m_handles[w], r->m_span, &r->m_prop, 0) != hipSuccess) + return fail("could not allocate " + std::to_string(r->m_span) + + " bytes of device memory for placement", + false); + r->m_created[w] = true; + if(hipMemMap(at, r->m_span, 0, r->m_handles[w], 0) != hipSuccess) + return fail("could not map placed memory", false); + r->m_mapped[w] = true; + if(hipMemSetAccess(at, r->m_span, &access, 1) != hipSuccess) + return fail("could not enable access to placed memory", false); + } + if(r->fill_poison(nullptr) != hipSuccess || hipDeviceSynchronize() != hipSuccess) + return fail("could not fill the placement poison windows", false); + return r; +#endif + } + + ~PlacedRegion() + { + (void)hipDeviceSynchronize(); + for(int w = 0; w < 3; w++) + { + if(m_mapped[w]) + (void)hipMemUnmap(reinterpret_cast(m_windows[w]), m_span); + if(m_created[w]) + (void)hipMemRelease(m_handles[w]); + } + if(m_reserved) + (void)hipMemAddressFree(m_reserved, m_reserved_bytes); + } + + PlacedRegion(const PlacedRegion&) = delete; + PlacedRegion& operator=(const PlacedRegion&) = delete; + + void* ptr() const + { + return reinterpret_cast(m_windows[0]); + } + + // The first byte at or above the 4 GiB boundary the buffer crosses. + uint64_t boundary() const + { + return (m_windows[0] / k4GiB + 1) * k4GiB; + } + + // Refills both poison windows with kFastCheckPoisonValue. + hipError_t fill_poison(hipStream_t stream) const + { + const size_t elements = m_span / fast_check_element_size(m_type); + for(int w = 1; w < 3; w++) + fast_check_poison_padding_device( + {reinterpret_cast(m_windows[w]), m_type, 0, 0, 1, 0}, 1, elements, stream); + return hipGetLastError(); + } + + // Fails when any element of either poison window no longer holds the poison value: a write + // reached an address exactly 4 GiB away from where it belonged. + FastCheckResult verify_poison(const std::string& operand, hipStream_t stream) const + { + FastCheckResult result; + std::ostringstream msg; + const size_t es = fast_check_element_size(m_type); + const char* where[3] = {"", "below", "above"}; + for(int w = 1; w < 3; w++) + { + FastCheckChanged c = fast_check_count_changed_device( + reinterpret_cast(m_windows[w]), m_type, m_span / es, stream); + if(!c.ok) + return {false, "could not scan the placement poison windows"}; + if(c.count) + { + result.passed = false; + msg << c.count << " elements of the poison window 4 GiB " << where[w] << " " + << operand << " were written. The first is " << operand << " element " + << c.first << " (byte offset " << c.first * es << ") moved 4 GiB " << where[w] + << ", the effect of an address that " + << (w == 1 ? "lost the carry out of its low 32 bits" + : "gained a carry into its high 32 bits") + << ".\n"; + } + } + result.message = msg.str(); + return result; + } + +private: + PlacedRegion() = default; + + hipDataType m_type{}; + hipMemAllocationProp m_prop{}; + size_t m_bytes = 0; + size_t m_span = 0; // bytes mapped per window + void* m_reserved = nullptr; + size_t m_reserved_bytes = 0; + uint64_t m_windows[3] = {}; // buffer, poison below, poison above + hipMemGenericAllocationHandle_t m_handles[3]{}; + bool m_created[3] = {}; + bool m_mapped[3] = {}; +}; diff --git a/projects/hipblaslt/clients/common/include/hipblaslt_arguments.hpp b/projects/hipblaslt/clients/common/include/hipblaslt_arguments.hpp index 2c212852477d..8c4ffb1f3284 100644 --- a/projects/hipblaslt/clients/common/include/hipblaslt_arguments.hpp +++ b/projects/hipblaslt/clients/common/include/hipblaslt_arguments.hpp @@ -174,6 +174,10 @@ struct Arguments int8_t unit_check; int8_t ulp_check; int8_t fast_check; + // fast_check: operand to place across a 4 GiB boundary (a, b, c, d, bias, + // scale_alpha_vec, workspace), and the bytes before the boundary (0 = middle) + char placement[16]; + size_t placement_offset; int8_t timing; char transA; @@ -305,6 +309,8 @@ struct Arguments OPER(unit_check) SEP \ OPER(ulp_check) SEP \ OPER(fast_check) SEP \ + OPER(placement) SEP \ + OPER(placement_offset) SEP \ OPER(timing) SEP \ OPER(transA) SEP \ OPER(transB) SEP \ diff --git a/projects/hipblaslt/clients/common/include/testing_matmul.hpp b/projects/hipblaslt/clients/common/include/testing_matmul.hpp index f8979ecac52f..30a5d7a8e43a 100644 --- a/projects/hipblaslt/clients/common/include/testing_matmul.hpp +++ b/projects/hipblaslt/clients/common/include/testing_matmul.hpp @@ -1131,6 +1131,32 @@ inline std::string fast_check_unsupported_reason(const Arguments& arg, return {}; } +// Skips the test when buffer placement is unavailable on this platform, and fails it when the +// placement request is invalid. +#ifdef GOOGLE_TEST +#define CHECK_PLACEMENT(ok, unsupported, why) \ + do \ + { \ + if(!(ok)) \ + { \ + if(unsupported) \ + GTEST_SKIP() << (why); \ + else \ + FAIL() << (why); \ + } \ + } while(0) +#else +#define CHECK_PLACEMENT(ok, unsupported, why) \ + do \ + { \ + if(!(ok)) \ + { \ + hipblaslt_cerr << (why) << std::endl; \ + return; \ + } \ + } while(0) +#endif + // Seed for the fast_check probe vectors: FNV-1a over the test name, so a failure reproduces. inline uint64_t fast_check_seed(const Arguments& arg) { @@ -2368,6 +2394,36 @@ void testing_matmul_with_bias(const Arguments& arg, } } + if(arg.placement[0]) + { + static const char* operands[] + = {"a", "b", "c", "d", "bias", "scale_alpha_vec", "workspace"}; + std::string why; + if(!arg.fast_check) + why = "placement requires fast_check"; + else if(std::none_of(std::begin(operands), std::end(operands), [&](const char* o) { + return !strcmp(o, arg.placement); + })) + why = std::string("unknown placement operand '") + arg.placement + "'"; + else if(HMM) + why = "placement does not support HMM"; + else if(arg.c_equal_d && !strcmp(arg.placement, "d")) + why = "with c_equal_d, D is C: place c instead of d"; + else if(!strcmp(arg.placement, "bias") && !arg.bias_vector) + why = "placing the bias requires bias_vector"; + else if(!strcmp(arg.placement, "scale_alpha_vec") && !arg.scaleAlpha_vector) + why = "placing the scaleAlpha vector requires scaleAlpha_vector"; + if(!why.empty()) + { +#ifdef GOOGLE_TEST + FAIL() << why; +#else + hipblaslt_cerr << why << std::endl; + return; +#endif + } + } + // Calculating block count auto plan = hipblaslt_bench::compute_rotating_buffer_plan( arg.adaptive, arg.max_iters, arg.cold_iters, arg.iters, rotating, totalRotatingSizeNeeded); @@ -2665,17 +2721,44 @@ void testing_matmul_with_bias(const Arguments& arg, epilogue_on[i] = true; } - // allocate memory on device - dA.emplace_back(TiA, size_dA[i] * block_count, HMM); + // allocate memory on device; the operand named by arg.placement crosses a 4 GiB + // boundary. allocate() returns false when that placement failed. + std::string placement_why; + bool placement_unsupported = false; + auto allocate = [&](std::vector& v, + hipDataType type, + size_t elements, + const char* operand) { + if(i == 0 && !strcmp(arg.placement, operand)) + { + v.emplace_back(type, + elements, + size_t(arg.placement_offset), + &placement_why, + &placement_unsupported); + return v.back().buf() != nullptr; + } + v.emplace_back(type, elements, HMM); + return true; + }; + CHECK_PLACEMENT(allocate(dA, TiA, size_dA[i] * block_count, "a"), + placement_unsupported, + placement_why); CHECK_DEVICE_ALLOCATION(hipGetLastError()); - dB.emplace_back(TiB, size_dB[i] * block_count, HMM); + CHECK_PLACEMENT(allocate(dB, TiB, size_dB[i] * block_count, "b"), + placement_unsupported, + placement_why); CHECK_DEVICE_ALLOCATION(hipGetLastError()); - dC.emplace_back(To, size_C[i] * block_count, HMM); + CHECK_PLACEMENT(allocate(dC, To, size_C[i] * block_count, "c"), + placement_unsupported, + placement_why); CHECK_DEVICE_ALLOCATION(hipGetLastError()); if(!arg.c_equal_d) { - dD.emplace_back(To, size_D[i] * block_count, HMM); + CHECK_PLACEMENT(allocate(dD, To, size_D[i] * block_count, "d"), + placement_unsupported, + placement_why); CHECK_DEVICE_ALLOCATION(hipGetLastError()); dDp = &dD; } @@ -2684,13 +2767,20 @@ void testing_matmul_with_bias(const Arguments& arg, if(size_bias[i] * block_count != 0) { - dBias.emplace_back(Tbias, size_bias[i] * block_count, HMM); + CHECK_PLACEMENT(allocate(dBias, Tbias, size_bias[i] * block_count, "bias"), + placement_unsupported, + placement_why); CHECK_DEVICE_ALLOCATION(hipGetLastError()); } if(arg.scaleAlpha_vector) { - dScaleAlphaVec.emplace_back(Talpha, size_scaleAlphaVec[i] * block_count, HMM); + CHECK_PLACEMENT(allocate(dScaleAlphaVec, + Talpha, + size_scaleAlphaVec[i] * block_count, + "scale_alpha_vec"), + placement_unsupported, + placement_why); CHECK_DEVICE_ALLOCATION(hipGetLastError()); } @@ -5021,6 +5111,22 @@ void testing_matmul_with_bias(const Arguments& arg, dWorkspace = new device_vector(workspace_size * block_count, 1, HMM); CHECK_DEVICE_ALLOCATION(dWorkspace->memcheck()); + // The workspace the checked solutions use. A placed workspace replaces the normal one. + std::unique_ptr placedWorkspace; + void* workspacePtr = static_cast(*dWorkspace); + if(!strcmp(arg.placement, "workspace")) + { + std::string why; + bool unsupported = false; + if(workspace_size == 0) + why = "placing the workspace requires a nonzero workspace size"; + else + placedWorkspace = PlacedRegion::create( + workspace_size, size_t(arg.placement_offset), HIP_R_8I, &why, &unsupported); + CHECK_PLACEMENT(placedWorkspace != nullptr, unsupported, why); + workspacePtr = placedWorkspace->ptr(); + } + if(arg.use_user_args) { CHECK_HIP_ERROR( @@ -5509,7 +5615,7 @@ void testing_matmul_with_bias(const Arguments& arg, CHECK_HIPBLASLT_ERROR( gemmVec[0].initialize(heuristicResult[sol].algo, tuningVec[heuristicTuningIndex[sol]], - *dWorkspace)); + workspacePtr)); CHECK_HIPBLASLT_ERROR(gemmVec[0].run(stream)); } else if(batchMode == HIPBLASLT_BATCH_MODE_POINTER_ARRAY) //For General Batch GEMM @@ -5537,7 +5643,7 @@ void testing_matmul_with_bias(const Arguments& arg, ddd[0], matD[0], &heuristicResult[sol].algo, - *dWorkspace, + workspacePtr, workspace_size, stream), HIPBLAS_STATUS_SUCCESS); @@ -5558,7 +5664,7 @@ void testing_matmul_with_bias(const Arguments& arg, (*dDp)[0].buf(), matD[0], &heuristicResult[sol].algo, - *dWorkspace, + workspacePtr, workspace_size, stream), HIPBLAS_STATUS_SUCCESS); @@ -5573,7 +5679,7 @@ void testing_matmul_with_bias(const Arguments& arg, CHECK_HIPBLASLT_ERROR( groupedGemmVec[0].initialize(heuristicResult[sol].algo, tuningVec[heuristicTuningIndex[0]], - *dWorkspace)); + workspacePtr)); groupedGemmVec[0].getDefaultValueForDeviceUserArguments(userArgs); // Copy them to device memory CHECK_HIP_ERROR(hipMemcpy(d_userArgs, @@ -5589,7 +5695,7 @@ void testing_matmul_with_bias(const Arguments& arg, CHECK_HIPBLASLT_ERROR( groupedGemmVec[0].initialize(heuristicResult[sol].algo, tuningVec[heuristicTuningIndex[0]], - *dWorkspace, + workspacePtr, false, stream)); @@ -5718,7 +5824,7 @@ void testing_matmul_with_bias(const Arguments& arg, {"B", dB[i].buf(), size_B[i] * realDataTypeSize(TiB)}, {"C", dC[i].buf(), size_C[i] * realDataTypeSize(To)}, {"D", (*dDp)[i].buf(), size_D[i] * realDataTypeSize(To)}, - {"workspace", static_cast(*dWorkspace), workspace_size}}; + {"workspace", workspacePtr, workspace_size}}; if(arg.bias_vector) buffers.push_back( {"bias", dBias[i].buf(), size_bias[i] * realDataTypeSize(Tbias)}); @@ -5742,6 +5848,33 @@ void testing_matmul_with_bias(const Arguments& arg, #endif } } + + // A write that missed a placed operand by exactly 4 GiB lands in its poison. + const PlacedRegion* placed = placedWorkspace.get(); + for(auto* v : {&dA, &dB, &dC, &dD, &dBias, &dScaleAlphaVec}) + if(!placed && !v->empty() && (*v)[0].placement()) + placed = (*v)[0].placement(); + if(placed) + { + FastCheckResult poison = placed->verify_poison(arg.placement, stream); + if(!poison.passed) + { + const std::string report + = "fast_check placement, " + + solution_description(handle, + heuristicResult[sol].algo, + sol, + heuristicResult.size(), + true) + + ":\n" + poison.message; +#ifdef GOOGLE_TEST + ADD_FAILURE() << report; +#else + hipblaslt_cerr << report << std::endl; +#endif + } + CHECK_HIP_ERROR(placed->fill_poison(stream)); + } } } } diff --git a/projects/hipblaslt/clients/common/src/fast_check.cpp b/projects/hipblaslt/clients/common/src/fast_check.cpp index 4c1a19e4cd38..547ba5a6584d 100644 --- a/projects/hipblaslt/clients/common/src/fast_check.cpp +++ b/projects/hipblaslt/clients/common/src/fast_check.cpp @@ -1447,48 +1447,97 @@ hipError_t fast_check_poison_padding_device(const FastCheckMatrix& m, return fill_outside_region(m, batch_count, total_elements, poison_bits(m.type), stream); } +namespace +{ + // Runs scan_kernel and returns its counters: [0] padding mismatches, [1] first padding + // mismatch, [2] unwritten elements, [3] first unwritten element. + hipError_t run_scan(const FastCheckMatrix& m, + int64_t batch_count, + size_t total_elements, + bool expect_poison, + hipStream_t stream, + unsigned long long (&counters)[4]) + { + unsigned long long init[4] = {0, ~0ull, 0, ~0ull}; + DeviceArray d_counters(4); + if(!d_counters.ptr) + return hipErrorOutOfMemory; + hipError_t err + = hipMemcpyAsync(d_counters.ptr, init, sizeof(init), hipMemcpyHostToDevice, stream); + if(err != hipSuccess) + return err; + // A failed launch, or an unsupported type, would otherwise leave the counters clean. + err = hipErrorInvalidValue; + switch(element_size(m.type)) + { + case 1: + err = launch_scan( + m, batch_count, total_elements, expect_poison, d_counters.ptr, stream); + break; + case 2: + err = launch_scan( + m, batch_count, total_elements, expect_poison, d_counters.ptr, stream); + break; + case 4: + err = launch_scan( + m, batch_count, total_elements, expect_poison, d_counters.ptr, stream); + break; + case 8: + err = launch_scan( + m, batch_count, total_elements, expect_poison, d_counters.ptr, stream); + break; + default: + break; + } + if(err != hipSuccess) + return err; + err = hipMemcpyAsync( + counters, d_counters.ptr, sizeof(counters), hipMemcpyDeviceToHost, stream); + if(err == hipSuccess) + err = hipStreamSynchronize(stream); + return err; + } +} // namespace + +size_t fast_check_element_size(hipDataType type) +{ + return element_size(type); +} + +uint64_t fast_check_poison_bits(hipDataType type) +{ + return poison_bits(type); +} + +FastCheckChanged fast_check_count_changed_device(const void* data, + hipDataType type, + size_t elements, + hipStream_t stream) +{ + // An empty region makes every element padding, so the scan compares all of them with the + // poison value. + FastCheckMatrix m{data, type, 0, 0, 1, 0}; + unsigned long long counters[4]; + FastCheckChanged changed; + if(run_scan(m, 1, elements, true, stream, counters) != hipSuccess) + { + changed.ok = false; + return changed; + } + changed.count = counters[0]; + changed.first = counters[1]; + return changed; +} + FastCheckResult fast_check_scan_padding_device(const FastCheckMatrix& m, int64_t batch_count, size_t total_elements, bool expect_poison, hipStream_t stream) { - FastCheckResult result; - unsigned long long init[4] = {0, ~0ull, 0, ~0ull}; - DeviceArray d_counters(4); - if(!d_counters.ptr - || hipMemcpyAsync(d_counters.ptr, init, sizeof(init), hipMemcpyHostToDevice, stream) - != hipSuccess) - return {false, "fast_check could not allocate its scan counters"}; - // A failed launch, or an unsupported type, would otherwise leave the counters clean. - hipError_t err = hipErrorInvalidValue; - switch(element_size(m.type)) - { - case 1: - err = launch_scan( - m, batch_count, total_elements, expect_poison, d_counters.ptr, stream); - break; - case 2: - err = launch_scan( - m, batch_count, total_elements, expect_poison, d_counters.ptr, stream); - break; - case 4: - err = launch_scan( - m, batch_count, total_elements, expect_poison, d_counters.ptr, stream); - break; - case 8: - err = launch_scan( - m, batch_count, total_elements, expect_poison, d_counters.ptr, stream); - break; - default: - break; - } - if(err != hipSuccess) - return {false, std::string("fast_check scan failed: ") + hipGetErrorString(err)}; + FastCheckResult result; unsigned long long counters[4]; - err = hipMemcpyAsync(counters, d_counters.ptr, sizeof(counters), hipMemcpyDeviceToHost, stream); - if(err == hipSuccess) - err = hipStreamSynchronize(stream); + hipError_t err = run_scan(m, batch_count, total_elements, expect_poison, stream, counters); if(err != hipSuccess) return {false, std::string("fast_check scan failed: ") + hipGetErrorString(err)}; diff --git a/projects/hipblaslt/clients/common/src/hipblaslt_arguments.cpp b/projects/hipblaslt/clients/common/src/hipblaslt_arguments.cpp index 56eb4c9a402f..10c842ef58c6 100644 --- a/projects/hipblaslt/clients/common/src/hipblaslt_arguments.cpp +++ b/projects/hipblaslt/clients/common/src/hipblaslt_arguments.cpp @@ -118,6 +118,8 @@ void Arguments::init() unit_check = 1; ulp_check = 0; fast_check = 0; + placement[0] = 0; + placement_offset = 0; timing = 0; transA = '*'; diff --git a/projects/hipblaslt/clients/tests/data/hipblaslt_common.yaml b/projects/hipblaslt/clients/tests/data/hipblaslt_common.yaml index 378d152498ec..ce2b7d0db3ae 100644 --- a/projects/hipblaslt/clients/tests/data/hipblaslt_common.yaml +++ b/projects/hipblaslt/clients/tests/data/hipblaslt_common.yaml @@ -648,6 +648,8 @@ Arguments: - unit_check: c_int8 - ulp_check: c_int8 - fast_check: c_int8 + - placement: c_char*16 + - placement_offset: c_size_t - timing: c_int8 - transA: c_char - transB: c_char @@ -757,6 +759,8 @@ Defaults: unit_check: 1 ulp_check: 0 fast_check: 0 + placement: '' + placement_offset: 0 timing: 0 iters: 10 cold_iters: 2 diff --git a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml index 21b47ca5b93f..6dc28808041d 100755 --- a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml +++ b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml @@ -3384,6 +3384,111 @@ Tests: gpu_arch: '9(0a|42|50)' +############################################################################### +# fast_check placement (AIHPBLAS-4994). +# +# Some address defects depend on where a buffer sits, not on problem size: a +# kernel that drops the carry out of the low 32 bits of an address (ROCM-32046) +# goes wrong only when a buffer crosses a 4 GiB boundary, and the wrong access +# lands exactly 4 GiB away, outside every allocation. These cases place one +# operand at a time across a 4 GiB virtual-address boundary (by default half of +# it on each side) with poison memory mapped 4 GiB below and above it. A +# dropped-carry read returns poison and fast_check reports a wrong D; a +# dropped-carry write changes the poison and the test reports it. The defect +# does not depend on size, so small shapes suffice, and every solution runs +# because such defects are specific to a few kernels. +############################################################################### + +# Every PR: each operand placed in turn, with the heuristic's first 16 solutions. +# The workspace is placed only in the all-solutions cases below: at this shape +# none of the first 16 f32 solutions uses one. +# Runtime on gfx942: 8 cases, about 4 s. +- name: matmul_fast_check_placement + category: pre_checkin + function: + matmul: + - *single_precision + - *hpa_bf16_precision + matrix_size: + - { M: 1031, N: 517, K: 263 } + transA: N + transB: N + alpha: 2.0 + beta: -2.0 + placement: [ a, b, c, d ] + requested_solution_num: 16 + initialization: integer_exact + unit_check: 0 + fast_check: 1 + +# The side operands, with the epilogue that reads them. +# Runtime on gfx942: 4 cases, about 4 s. +- name: matmul_fast_check_placement_epilogue + category: pre_checkin + function: + matmul: + - *single_precision + - *hpa_bf16_precision + matrix_size: + - { M: 1031, N: 517, K: 263 } + transA: N + transB: T + alpha: 1.0 + beta: 0.0 + bias_vector: 1 + bias_type: f32_r + scaleAlpha_vector: 1 + placement: [ bias, scale_alpha_vec ] + requested_solution_num: 16 + initialization: integer_exact + unit_check: 0 + fast_check: 1 + +# Nightly: the same placements with every supported solution, about 6 minutes on +# gfx942. +# Runtime on gfx942: 10 cases, about 178 s. +- name: matmul_fast_check_placement_all_solutions + category: nightly + function: + matmul: + - *single_precision + - *hpa_bf16_precision + matrix_size: + - { M: 1031, N: 517, K: 263 } + transA: N + transB: N + alpha: 2.0 + beta: -2.0 + placement: [ a, b, c, d, workspace ] + algo_method: [1] + requested_solution_num: -1 + initialization: integer_exact + unit_check: 0 + fast_check: 1 + +# Runtime on gfx942: 4 cases, about 190 s. +- name: matmul_fast_check_placement_epilogue_all_solutions + category: nightly + function: + matmul: + - *single_precision + - *hpa_bf16_precision + matrix_size: + - { M: 1031, N: 517, K: 263 } + transA: N + transB: T + alpha: 1.0 + beta: 0.0 + bias_vector: 1 + bias_type: f32_r + scaleAlpha_vector: 1 + placement: [ bias, scale_alpha_vec ] + algo_method: [1] + requested_solution_num: -1 + initialization: integer_exact + unit_check: 0 + fast_check: 1 + - name: matmul_96_complex category: pre_checkin function: diff --git a/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp b/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp index a1a4c39ff027..8af43863efc2 100644 --- a/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp +++ b/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp @@ -12,6 +12,8 @@ #include #include "fast_check.hpp" +#include "hipBuffer.hpp" +#include "hip_placement.hpp" #include @@ -1067,6 +1069,125 @@ namespace EXPECT_FALSE(res.message.empty()); } + // Reads word `word` of `base` twice: through the correct address, and through an address whose + // low 32 bits were computed without carrying into the high 32 bits. + __global__ void read_with_dropped_carry(const uint32_t* base, + uint64_t word, + uint32_t* correct, + uint32_t* dropped) + { + uint64_t b = reinterpret_cast(base); + uint64_t right = b + word * 4; + uint64_t wrapped = (b & ~uint64_t(0xffffffff)) | uint32_t(uint32_t(b) + uint32_t(word * 4)); + *correct = *reinterpret_cast(right); + *dropped = *reinterpret_cast(wrapped); + } + + __global__ void write_with_dropped_carry(uint32_t* base, uint64_t word, uint32_t value) + { + uint64_t b = reinterpret_cast(base); + uint64_t wrapped = (b & ~uint64_t(0xffffffff)) | uint32_t(uint32_t(b) + uint32_t(word * 4)); + *reinterpret_cast(wrapped) = value; + } + + // The placement harness must turn a dropped-carry read into a poison value instead of a fault, + // and report a dropped-carry write with the element it was meant for. + TEST(FastCheckDevice_pre_checkin, placed_region_catches_dropped_carry) + { + const size_t bytes = 8 << 20, below = 4 << 20; + std::string why; + bool unsupported = false; + auto region = PlacedRegion::create(bytes, below, HIP_R_32F, &why, &unsupported); + if(!region && unsupported) + GTEST_SKIP() << why; + ASSERT_TRUE(region) << why; + + const uint64_t start = reinterpret_cast(region->ptr()); + ASSERT_EQ(region->boundary() - start, below); + + std::vector data(bytes / 4); + for(size_t n = 0; n < data.size(); n++) + data[n] = uint32_t(n); + ASSERT_EQ(hipMemcpy(region->ptr(), data.data(), bytes, hipMemcpyHostToDevice), hipSuccess); + + const uint64_t word = below / 4 + 100; // 100 words past the boundary + uint32_t* d_out; + ASSERT_EQ(hipMalloc(&d_out, 2 * sizeof(uint32_t)), hipSuccess); + hipLaunchKernelGGL(read_with_dropped_carry, + dim3(1), + dim3(1), + 0, + 0, + static_cast(region->ptr()), + word, + d_out, + d_out + 1); + uint32_t out[2]; + ASSERT_EQ(hipMemcpy(out, d_out, sizeof(out), hipMemcpyDeviceToHost), hipSuccess); + (void)hipFree(d_out); + float poison = kFastCheckPoisonValue; + EXPECT_EQ(out[0], uint32_t(word)); + EXPECT_EQ(out[1], *reinterpret_cast(&poison)); + + auto clean = region->verify_poison("A", 0); + EXPECT_TRUE(clean.passed) << clean.message; + + hipLaunchKernelGGL(write_with_dropped_carry, + dim3(1), + dim3(1), + 0, + 0, + static_cast(region->ptr()), + word, + 0x12345678u); + ASSERT_EQ(hipDeviceSynchronize(), hipSuccess); + auto written = region->verify_poison("A", 0); + ASSERT_FALSE(written.passed); + EXPECT_NE( + written.message.find("1 elements of the poison window 4 GiB below A were written"), + std::string::npos) + << written.message; + EXPECT_NE(written.message.find("A element " + std::to_string(word)), std::string::npos) + << written.message; + + ASSERT_EQ(region->fill_poison(0), hipSuccess); + auto refilled = region->verify_poison("A", 0); + EXPECT_TRUE(refilled.passed) << refilled.message; + } + + TEST(FastCheckDevice_pre_checkin, placement_rejects_requests_that_cannot_cross) + { + std::string why; + bool unsupported = false; + if(!PlacedRegion::create(1 << 20, 0, HIP_R_32F, &why, &unsupported) && unsupported) + GTEST_SKIP() << why; + EXPECT_FALSE(PlacedRegion::create(4096, 8192, HIP_R_32F, &why, &unsupported)); + EXPECT_FALSE(unsupported); + EXPECT_NE(why.find("cannot cross"), std::string::npos) << why; + EXPECT_FALSE(PlacedRegion::create(1 << 20, 100, HIP_R_32F, &why, &unsupported)); + EXPECT_NE(why.find("multiple of the mapping granularity"), std::string::npos) << why; + } + + // std::vector growth moves its elements. A placed buffer must keep its address through the + // move, and the moved-from object must not touch the placed memory when it is destroyed. + TEST(FastCheckDevice_pre_checkin, placed_buffer_survives_vector_growth) + { + std::string why; + bool unsupported = false; + std::vector buffers; + buffers.emplace_back(HIP_R_32F, size_t(1 << 20), size_t(0), &why, &unsupported); + if(!buffers.back().buf() && unsupported) + GTEST_SKIP() << why; + ASSERT_TRUE(buffers.back().buf()) << why; + void* placed = buffers[0].buf(); + for(int n = 0; n < 16; n++) + buffers.emplace_back(HIP_R_32F, size_t(1024)); + EXPECT_EQ(buffers[0].buf(), placed); + ASSERT_TRUE(buffers[0].placement()); + EXPECT_EQ(buffers[0].placement()->ptr(), placed); + EXPECT_EQ(hipMemset(placed, 0, 4 << 20), hipSuccess); + } + TEST(FastCheckDevice_pre_checkin, copy_region_to_host_drops_the_padding) { DeviceMatrix m; diff --git a/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp b/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp index 86f234388aef..c7bcbac8cffc 100644 --- a/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp +++ b/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp @@ -171,6 +171,8 @@ namespace if(arg.c_equal_d) name << "_C_EQUAL_D"; + if(arg.placement[0]) + name << "_PLACE_" << arg.placement; // grouped gemm only supports ext if(arg.use_ext || arg.grouped_gemm > 0) name << "_APIExt"; From 80d51f945cd51ed09ae916d35b821c81163b2b4c Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Sat, 3 Oct 2026 10:29:06 -0400 Subject: [PATCH 2/8] test(hipblaslt): AIHPBLAS-4994 skip the normal workspace when the workspace is placed A placement run that places the workspace still allocated the full normal workspace, which nothing used: placement requires fast_check, fast_check refuses timing runs, and the checked solutions use the placed workspace. For a large workspace that extra allocation could run out of memory before the placed one. The normal workspace is now empty in that case. Co-authored-by: Cursor --- .../clients/common/include/testing_matmul.hpp | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/projects/hipblaslt/clients/common/include/testing_matmul.hpp b/projects/hipblaslt/clients/common/include/testing_matmul.hpp index 30a5d7a8e43a..482889cca206 100644 --- a/projects/hipblaslt/clients/common/include/testing_matmul.hpp +++ b/projects/hipblaslt/clients/common/include/testing_matmul.hpp @@ -5108,13 +5108,17 @@ void testing_matmul_with_bias(const Arguments& arg, CHECK_SOLUTION_FOUND(returnedAlgoCount); - dWorkspace = new device_vector(workspace_size * block_count, 1, HMM); + // A placed workspace replaces the normal one, which placement runs never use: placement + // requires fast_check, and fast_check refuses timing runs. + const bool placeWorkspace = !strcmp(arg.placement, "workspace"); + dWorkspace = new device_vector( + placeWorkspace ? 0 : workspace_size * block_count, 1, HMM); CHECK_DEVICE_ALLOCATION(dWorkspace->memcheck()); - // The workspace the checked solutions use. A placed workspace replaces the normal one. + // The workspace the checked solutions use. std::unique_ptr placedWorkspace; void* workspacePtr = static_cast(*dWorkspace); - if(!strcmp(arg.placement, "workspace")) + if(placeWorkspace) { std::string why; bool unsupported = false; From 284b46d2fa21a9adc0e40790de5db11cb0d5aa6c Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Sat, 3 Oct 2026 10:29:12 -0400 Subject: [PATCH 3/8] test(hipblaslt): AIHPBLAS-4994 cover the poison helpers for every type fast_check_poison_bits and fast_check_count_changed_device had no self-test, and Codecov reported 69% patch coverage for this layer. A new self-test checks the poison bits and the changed-element count, with its first index, in every supported type. Co-authored-by: Cursor --- .../clients/tests/src/fast_check_gtest.cpp | 32 +++++++++++++++++++ 1 file changed, 32 insertions(+) diff --git a/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp b/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp index 8af43863efc2..e7fe5d256069 100644 --- a/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp +++ b/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp @@ -927,6 +927,38 @@ namespace } } + // In every supported type, the poison bits are kFastCheckPoisonValue as the type stores it, + // and the count of changed poison finds a single changed element. + TEST(FastCheckDevice_pre_checkin, every_type_counts_changed_poison) + { + const size_t elements = 300; + for(const auto& tc : kTypes) + { + std::vector poison(tc.size); + store_as(tc.type, kFastCheckPoisonValue, poison.data()); + uint64_t bits = fast_check_poison_bits(tc.type); + EXPECT_EQ(std::memcmp(&bits, poison.data(), tc.size), 0) << tc.name; + + std::vector h(elements * tc.size); + for(size_t idx = 0; idx < elements; idx++) + std::memcpy(h.data() + idx * tc.size, poison.data(), tc.size); + char* d = nullptr; + ASSERT_EQ(hipMalloc(&d, h.size()), hipSuccess); + ASSERT_EQ(hipMemcpy(d, h.data(), h.size(), hipMemcpyHostToDevice), hipSuccess); + auto changed = fast_check_count_changed_device(d, tc.type, elements, 0); + EXPECT_TRUE(changed.ok) << tc.name; + EXPECT_EQ(changed.count, 0u) << tc.name; + + store_as(tc.type, 1, h.data() + 217 * tc.size); + ASSERT_EQ(hipMemcpy(d, h.data(), h.size(), hipMemcpyHostToDevice), hipSuccess); + changed = fast_check_count_changed_device(d, tc.type, elements, 0); + (void)hipFree(d); + EXPECT_TRUE(changed.ok) << tc.name; + EXPECT_EQ(changed.count, 1u) << tc.name; + EXPECT_EQ(changed.first, 217u) << tc.name; + } + } + // Configurations fast_check cannot check exactly are refused with a reason, before any sums. TEST(FastCheck_pre_checkin, unsupported_configurations_are_refused) { From cfbea3bcbf3555b2fc631bb625e361b29c074427 Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Sat, 3 Oct 2026 10:48:13 -0400 Subject: [PATCH 4/8] test(hipblaslt): AIHPBLAS-4994 skip workspace placement when no solution uses a workspace On gfx12 every solution of the all-solutions workspace-placement cases needs no workspace, and the case failed with "placing the workspace requires a nonzero workspace size" in the Math CI precheckin run. Whether a workspace is needed depends on the library and GPU, so that case now skips. Co-authored-by: Cursor --- .../hipblaslt/clients/common/include/testing_matmul.hpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/projects/hipblaslt/clients/common/include/testing_matmul.hpp b/projects/hipblaslt/clients/common/include/testing_matmul.hpp index 482889cca206..dbcb45396aa2 100644 --- a/projects/hipblaslt/clients/common/include/testing_matmul.hpp +++ b/projects/hipblaslt/clients/common/include/testing_matmul.hpp @@ -5122,8 +5122,13 @@ void testing_matmul_with_bias(const Arguments& arg, { std::string why; bool unsupported = false; + // Whether any solution needs a workspace depends on the library and GPU, so a run with + // none has nothing to place: a skip, not a failure. if(workspace_size == 0) - why = "placing the workspace requires a nonzero workspace size"; + { + why = "no solution here uses a workspace, so there is nothing to place"; + unsupported = true; + } else placedWorkspace = PlacedRegion::create( workspace_size, size_t(arg.placement_offset), HIP_R_8I, &why, &unsupported); From 022648bf96a1f561893bd506d15bfdfefd07a2c2 Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Sat, 3 Oct 2026 12:19:03 -0400 Subject: [PATCH 5/8] test(hipblaslt): AIHPBLAS-4994 read the poison bits with memcpy in the placement self-test The self-test read a float's bits through a uint32_t pointer, which breaks strict aliasing and can be miscompiled. It now copies them with memcpy. Co-authored-by: Cursor --- projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp b/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp index e7fe5d256069..162af927d369 100644 --- a/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp +++ b/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp @@ -1157,9 +1157,11 @@ namespace uint32_t out[2]; ASSERT_EQ(hipMemcpy(out, d_out, sizeof(out), hipMemcpyDeviceToHost), hipSuccess); (void)hipFree(d_out); - float poison = kFastCheckPoisonValue; + const float poison = kFastCheckPoisonValue; + uint32_t poison_bits; + std::memcpy(&poison_bits, &poison, sizeof(poison_bits)); EXPECT_EQ(out[0], uint32_t(word)); - EXPECT_EQ(out[1], *reinterpret_cast(&poison)); + EXPECT_EQ(out[1], poison_bits); auto clean = region->verify_poison("A", 0); EXPECT_TRUE(clean.passed) << clean.message; From 4e84c18305bad66801d36b15c6f721ee2f870400 Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Sat, 3 Oct 2026 15:38:46 -0400 Subject: [PATCH 6/8] test(hipblaslt): AIHPBLAS-4994 make the poison scans refuse an unsupported type run_scan launched nothing for an element size it does not handle and returned success with zero counters, so the padding scan and the changed-poison count reported a clean result for a type they never looked at. It now returns hipErrorInvalidValue, and a self-test checks both callers fail. Co-authored-by: Cursor --- projects/hipblaslt/clients/common/src/fast_check.cpp | 3 ++- .../hipblaslt/clients/tests/src/fast_check_gtest.cpp | 12 ++++++++++++ 2 files changed, 14 insertions(+), 1 deletion(-) diff --git a/projects/hipblaslt/clients/common/src/fast_check.cpp b/projects/hipblaslt/clients/common/src/fast_check.cpp index 547ba5a6584d..443b8756025e 100644 --- a/projects/hipblaslt/clients/common/src/fast_check.cpp +++ b/projects/hipblaslt/clients/common/src/fast_check.cpp @@ -1487,7 +1487,8 @@ namespace m, batch_count, total_elements, expect_poison, d_counters.ptr, stream); break; default: - break; + // An unsupported type would otherwise scan nothing and look clean. + return hipErrorInvalidValue; } if(err != hipSuccess) return err; diff --git a/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp b/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp index 162af927d369..c6eb1b38e866 100644 --- a/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp +++ b/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp @@ -959,6 +959,18 @@ namespace } } + // A type the scans do not support is an error, never a clean scan. + TEST(FastCheckDevice_pre_checkin, scans_refuse_an_unsupported_type) + { + DeviceMatrix m; + FastCheckMatrix complex = m.matrix(); + complex.type = HIP_C_32F; + auto res = fast_check_scan_padding_device(complex, DeviceMatrix::batch, 8, true, 0); + EXPECT_FALSE(res.passed); + EXPECT_NE(res.message.find("scan failed"), std::string::npos) << res.message; + EXPECT_FALSE(fast_check_count_changed_device(m.d, HIP_C_32F, 8, 0).ok); + } + // Configurations fast_check cannot check exactly are refused with a reason, before any sums. TEST(FastCheck_pre_checkin, unsupported_configurations_are_refused) { From 806c961ee42bdc609d595d1523f9772816d403fd Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Sun, 4 Oct 2026 16:43:26 -0400 Subject: [PATCH 7/8] fix(hipblaslt): AIHPBLAS-4994 make every solution's workspace straddle the boundary A placed workspace put the 4 GiB boundary at half the largest solution's workspace, so any solution needing less than half of it never reached the boundary and still passed as placed. The workspace region is now twice the largest size, and each solution starts so that its own reported range straddles the boundary while the whole workspace size, which any solution may use, stays mapped after it. The harness now fails a placement case whose buffer was not placed or whose range does not cross the boundary, instead of passing while testing nothing. The mapped tail of a placed buffer, after its last element, is poisoned and scanned like the windows 4 GiB away, and the poison fill reports launch errors. Co-authored-by: Cursor --- .../clients/common/include/hip_placement.hpp | 90 +++++++++++++++---- .../clients/common/include/testing_matmul.hpp | 51 ++++++++--- .../clients/tests/src/fast_check_gtest.cpp | 39 ++++++++ 3 files changed, 153 insertions(+), 27 deletions(-) diff --git a/projects/hipblaslt/clients/common/include/hip_placement.hpp b/projects/hipblaslt/clients/common/include/hip_placement.hpp index eda6fccaa4d9..ffb4c256b066 100644 --- a/projects/hipblaslt/clients/common/include/hip_placement.hpp +++ b/projects/hipblaslt/clients/common/include/hip_placement.hpp @@ -23,6 +23,8 @@ #include #include #include +#include +#include class PlacedRegion { @@ -156,14 +158,48 @@ class PlacedRegion return (m_windows[0] / k4GiB + 1) * k4GiB; } - // Refills both poison windows with kFastCheckPoisonValue. + // Bytes mapped for the buffer: its size rounded up to the mapping granularity. + size_t span() const + { + return m_span; + } + + // Whether [p, p + bytes) crosses the boundary. + bool crosses(const void* p, size_t bytes) const + { + const uint64_t at = reinterpret_cast(p); + return at < boundary() && at + bytes > boundary(); + } + + // A 256-byte aligned start inside the buffer window from which a range of `bytes` straddles + // the boundary near its middle, while `room` bytes from the start (at least `bytes`) stay + // inside the buffer. Different solutions use different amounts of one placed workspace, each + // needs its own range to cross, and each may touch up to the whole workspace size it was + // given. A range of 256 bytes or less cannot cross from an aligned start. + void* straddle(size_t bytes, size_t room) const + { + room = std::max(room, bytes); + if(bytes <= 256 || room > m_bytes) + return ptr(); + const uint64_t lo = m_windows[0]; + const uint64_t hi = (m_windows[0] + m_bytes - room) / 256 * 256; + const uint64_t half = std::max(256, bytes / 2 / 256 * 256); + return reinterpret_cast(std::clamp(boundary() - half, lo, std::max(lo, hi))); + } + + // Refills both poison windows, and the buffer window past the buffer's end, with + // kFastCheckPoisonValue. hipError_t fill_poison(hipStream_t stream) const { - const size_t elements = m_span / fast_check_element_size(m_type); - for(int w = 1; w < 3; w++) - fast_check_poison_padding_device( - {reinterpret_cast(m_windows[w]), m_type, 0, 0, 1, 0}, 1, elements, stream); - return hipGetLastError(); + hipError_t err = hipSuccess; + for(const auto& [at, elements] : poisoned_ranges()) + { + hipError_t e = fast_check_poison_padding_device( + {reinterpret_cast(at), m_type, 0, 0, 1, 0}, 1, elements, stream); + if(e != hipSuccess) + err = e; + } + return err; } // Fails when any element of either poison window no longer holds the poison value: a write @@ -172,25 +208,32 @@ class PlacedRegion { FastCheckResult result; std::ostringstream msg; - const size_t es = fast_check_element_size(m_type); - const char* where[3] = {"", "below", "above"}; - for(int w = 1; w < 3; w++) + const size_t es = fast_check_element_size(m_type); + const auto ranges = poisoned_ranges(); + for(size_t w = 0; w < ranges.size(); w++) { - FastCheckChanged c = fast_check_count_changed_device( - reinterpret_cast(m_windows[w]), m_type, m_span / es, stream); + const auto [at, elements] = ranges[w]; + FastCheckChanged c = fast_check_count_changed_device( + reinterpret_cast(at), m_type, elements, stream); if(!c.ok) return {false, "could not scan the placement poison windows"}; - if(c.count) + if(!c.count) + continue; + result.passed = false; + if(w < 2) { - result.passed = false; - msg << c.count << " elements of the poison window 4 GiB " << where[w] << " " - << operand << " were written. The first is " << operand << " element " - << c.first << " (byte offset " << c.first * es << ") moved 4 GiB " << where[w] + const char* where = w == 0 ? "below" : "above"; + msg << c.count << " elements of the poison window 4 GiB " << where << " " << operand + << " were written. The first is " << operand << " element " << c.first + << " (byte offset " << c.first * es << ") moved 4 GiB " << where << ", the effect of an address that " - << (w == 1 ? "lost the carry out of its low 32 bits" + << (w == 0 ? "lost the carry out of its low 32 bits" : "gained a carry into its high 32 bits") << ".\n"; } + else + msg << c.count << " elements past the end of " << operand + << " were written; the first is element " << m_bytes / es + c.first << ".\n"; } result.message = msg.str(); return result; @@ -199,6 +242,19 @@ class PlacedRegion private: PlacedRegion() = default; + // (address, elements) of each poisoned range: the windows 4 GiB below and above the buffer, + // then the mapped tail of the buffer window after the buffer's last element. + std::vector> poisoned_ranges() const + { + const size_t es = fast_check_element_size(m_type); + std::vector> ranges + = {{m_windows[1], m_span / es}, {m_windows[2], m_span / es}}; + const size_t used = (m_bytes + es - 1) / es; + if(m_span / es > used) + ranges.push_back({m_windows[0] + used * es, m_span / es - used}); + return ranges; + } + hipDataType m_type{}; hipMemAllocationProp m_prop{}; size_t m_bytes = 0; diff --git a/projects/hipblaslt/clients/common/include/testing_matmul.hpp b/projects/hipblaslt/clients/common/include/testing_matmul.hpp index dbcb45396aa2..355f1d1986b6 100644 --- a/projects/hipblaslt/clients/common/include/testing_matmul.hpp +++ b/projects/hipblaslt/clients/common/include/testing_matmul.hpp @@ -5115,9 +5115,11 @@ void testing_matmul_with_bias(const Arguments& arg, placeWorkspace ? 0 : workspace_size * block_count, 1, HMM); CHECK_DEVICE_ALLOCATION(dWorkspace->memcheck()); - // The workspace the checked solutions use. + // The workspace the checked solutions use. A placed workspace moves for each solution, so + // that solution's own workspace straddles the 4 GiB boundary. std::unique_ptr placedWorkspace; - void* workspacePtr = static_cast(*dWorkspace); + void* workspacePtr = static_cast(*dWorkspace); + size_t workspaceBytes = workspace_size; if(placeWorkspace) { std::string why; @@ -5130,8 +5132,14 @@ void testing_matmul_with_bias(const Arguments& arg, unsupported = true; } else - placedWorkspace = PlacedRegion::create( - workspace_size, size_t(arg.placement_offset), HIP_R_8I, &why, &unsupported); + // Twice the workspace (plus room for the boundary to land on a mapping granule), so + // that each solution's start can sit below the boundary with the whole workspace + // size, which any solution may use, still mapped after it. + placedWorkspace = PlacedRegion::create(2 * workspace_size + (size_t(4) << 20), + size_t(arg.placement_offset), + HIP_R_8I, + &why, + &unsupported); CHECK_PLACEMENT(placedWorkspace != nullptr, unsupported, why); workspacePtr = placedWorkspace->ptr(); } @@ -5592,6 +5600,12 @@ void testing_matmul_with_bias(const Arguments& arg, SCOPED_TRACE(solution_description( handle, heuristicResult[sol].algo, sol, heuristicResult.size(), false)); #endif + if(placedWorkspace) + { + workspacePtr + = placedWorkspace->straddle(heuristicResult[sol].workspaceSize, workspace_size); + workspaceBytes = heuristicResult[sol].workspaceSize; + } if(arg.fast_check && !arg.c_equal_d) { for(int i = 0; i < gemm_count; i++) @@ -5828,12 +5842,12 @@ void testing_matmul_with_bias(const Arguments& arg, FastCheckResult res = fast_check_result_device(fp, fcExpected[i], stream); if(!scan.passed || !res.passed) { - std::vector buffers = { - {"A", dA[i].buf(), size_A[i] * realDataTypeSize(TiA)}, - {"B", dB[i].buf(), size_B[i] * realDataTypeSize(TiB)}, - {"C", dC[i].buf(), size_C[i] * realDataTypeSize(To)}, - {"D", (*dDp)[i].buf(), size_D[i] * realDataTypeSize(To)}, - {"workspace", workspacePtr, workspace_size}}; + std::vector buffers + = {{"A", dA[i].buf(), size_A[i] * realDataTypeSize(TiA)}, + {"B", dB[i].buf(), size_B[i] * realDataTypeSize(TiB)}, + {"C", dC[i].buf(), size_C[i] * realDataTypeSize(To)}, + {"D", (*dDp)[i].buf(), size_D[i] * realDataTypeSize(To)}, + {"workspace", workspacePtr, workspaceBytes}}; if(arg.bias_vector) buffers.push_back( {"bias", dBias[i].buf(), size_bias[i] * realDataTypeSize(Tbias)}); @@ -5863,6 +5877,23 @@ void testing_matmul_with_bias(const Arguments& arg, for(auto* v : {&dA, &dB, &dC, &dD, &dBias, &dScaleAlphaVec}) if(!placed && !v->empty() && (*v)[0].placement()) placed = (*v)[0].placement(); + // A placement that silently fell back to a normal allocation would pass while + // testing nothing. + const bool straddles = placed + && (placed != placedWorkspace.get() || workspaceBytes <= 256 + || placed->crosses(workspacePtr, workspaceBytes)); + if(*arg.placement && !straddles) + { + const std::string report + = std::string("fast_check placement: ") + arg.placement + + " was to cross a 4 GiB boundary but " + + (placed ? "this solution's range does not" : "no buffer was placed"); +#ifdef GOOGLE_TEST + ADD_FAILURE() << report; +#else + hipblaslt_cerr << report << std::endl; +#endif + } if(placed) { FastCheckResult poison = placed->verify_poison(arg.placement, stream); diff --git a/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp b/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp index c6eb1b38e866..6235b2aa5319 100644 --- a/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp +++ b/projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp @@ -1214,6 +1214,45 @@ namespace EXPECT_NE(why.find("multiple of the mapping granularity"), std::string::npos) << why; } + // Solutions use different amounts of one placed workspace, and each one's range must cross the + // boundary while staying inside the mapped window. A write past the end of a placed buffer, + // into the rest of its mapped window, must be reported. + TEST(FastCheckDevice_pre_checkin, placed_workspace_straddles_for_every_size) + { + const size_t bytes = (size_t(16) << 20) - 4096; + std::string why; + bool unsupported = false; + auto region = PlacedRegion::create(bytes, 0, HIP_R_8I, &why, &unsupported); + if(!region && unsupported) + GTEST_SKIP() << why; + ASSERT_TRUE(region) << why; + + // As the harness places a workspace: twice the largest size, so any solution may use all + // of it from wherever its own range starts. + const uint64_t lo = reinterpret_cast(region->ptr()); + const size_t room = bytes / 2; + for(size_t ws : {size_t(300), size_t(4096), size_t(1) << 20, size_t(5) << 20, room}) + { + void* p = region->straddle(ws, room); + const uint64_t at = reinterpret_cast(p); + EXPECT_TRUE(region->crosses(p, ws)) << ws; + EXPECT_EQ(at % 256, 0u) << ws; + EXPECT_GE(at, lo) << ws; + EXPECT_LE(at + room, lo + bytes) << ws; + } + + auto clean = region->verify_poison("workspace", 0); + EXPECT_TRUE(clean.passed) << clean.message; + if(region->span() > bytes) + { + ASSERT_EQ(hipMemset(static_cast(region->ptr()) + bytes, 1, 1), hipSuccess); + auto tail = region->verify_poison("workspace", 0); + EXPECT_FALSE(tail.passed); + EXPECT_NE(tail.message.find("past the end of workspace"), std::string::npos) + << tail.message; + } + } + // std::vector growth moves its elements. A placed buffer must keep its address through the // move, and the moved-from object must not touch the placed memory when it is destroyed. TEST(FastCheckDevice_pre_checkin, placed_buffer_survives_vector_growth) From 4e3dcc99f913b373d3af59cafeed4ed1ce813a90 Mon Sep 17 00:00:00 2001 From: James Newling Date: Tue, 6 Oct 2026 11:53:22 -0700 Subject: [PATCH 8/8] fix(hipblaslt): release workspace on placement exits (cherry picked from commit 81db14636ea8f61a3291b59e0d0a5cd1d288c591) --- .../hipblaslt/clients/common/include/testing_matmul.hpp | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/projects/hipblaslt/clients/common/include/testing_matmul.hpp b/projects/hipblaslt/clients/common/include/testing_matmul.hpp index 355f1d1986b6..6a66c483a9fc 100644 --- a/projects/hipblaslt/clients/common/include/testing_matmul.hpp +++ b/projects/hipblaslt/clients/common/include/testing_matmul.hpp @@ -59,6 +59,7 @@ #include #include #include +#include #include #include #include @@ -4106,7 +4107,7 @@ void testing_matmul_with_bias(const Arguments& arg, HIPBLAS_STATUS_SUCCESS); // set workspace - device_vector* dWorkspace = nullptr; + std::unique_ptr> dWorkspace; size_t workspace_size = 0; // set user args @@ -5111,7 +5112,7 @@ void testing_matmul_with_bias(const Arguments& arg, // A placed workspace replaces the normal one, which placement runs never use: placement // requires fast_check, and fast_check refuses timing runs. const bool placeWorkspace = !strcmp(arg.placement, "workspace"); - dWorkspace = new device_vector( + dWorkspace = std::make_unique>( placeWorkspace ? 0 : workspace_size * block_count, 1, HMM); CHECK_DEVICE_ALLOCATION(dWorkspace->memcheck()); @@ -6706,8 +6707,6 @@ void testing_matmul_with_bias(const Arguments& arg, } } - if(dWorkspace != nullptr) - delete dWorkspace; if(userArgs != nullptr) CHECK_HIP_ERROR(hipFree(userArgs)); if(d_userArgs != nullptr)