leonarski_fandClaude Opus 5
373fda05a4
rc.172 stopped two CUDA failures from staying behind as the thread's last error, both inside
CUDAMemHelpers.h: a pooled allocation that fell back to cudaMalloc, and a stream that could not be
created. The reason those mattered was general, and four other places have it too.
Nothing but cudaGetLastError() takes an error away. A later successful call neither clears it nor
returns it - measured: after a cudaMalloc that cannot be served, cudaStreamSynchronize,
cudaMemcpyAsync and a small cudaMalloc all report success, and the "out of memory" is still waiting
for whoever calls cudaGetLastError() next. In this code that is the check after a kernel launch. So
wherever a CUDA failure is caught and the work carries on by another route, the error it left is
reported over the route that went fine:
- the device decoder in MXAnalysisWithoutFPGA::Analyze, whose failure is answered by decompressing
on the host - and the host route runs kernels of its own. This is the broker's path, and it is
the pair of log lines a long-running broker showed, one warning naming the device route and one
fatal error naming nothing.
- the indexer worker, which turns a failed attempt into a result and goes on to the next frame -
on the same thread, so one failure would read as an indexer that never works again.
- the beam stop's GPU accumulate and the FFT beam-centre capture, both of which fall back to the
host.
Each discharges it now, through cuda_clear_error() in CUDAWrapper (a no-op without CUDA). A sticky
error is not cleared by this and nothing here pretends otherwise.
tests/CudaHandledErrorTest.cpp: that a failed allocation leaves an error which survives later
successful work until it is cleared, and that an image whose device decode fails comes out of
Analyze whole and leaves nothing behind for the next one.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_013nW6FNRP1bBJJ8pfHiByAT
86 lines
3.7 KiB
C++
86 lines
3.7 KiB
C++
// SPDX-FileCopyrightText: 2026 Filip Leonarski, Paul Scherrer Institute
|
|
// SPDX-License-Identifier: GPL-3.0-only
|
|
|
|
#include <catch2/catch_all.hpp>
|
|
#include "../common/CUDAWrapper.h"
|
|
|
|
#ifdef JFJOCH_USE_CUDA
|
|
|
|
#include <cuda_runtime.h>
|
|
#include <limits>
|
|
|
|
#include "../common/AzimuthalIntegrationMapping.h"
|
|
#include "../common/AzimuthalIntegrationProfile.h"
|
|
#include "../common/DiffractionExperiment.h"
|
|
#include "../common/PixelMask.h"
|
|
#include "../compression/JFJochCompressor.h"
|
|
#include "../image_analysis/IndexAndRefine.h"
|
|
#include "../image_analysis/MXAnalysisWithoutFPGA.h"
|
|
#include "../image_analysis/indexing/CUDAMemHelpers.h"
|
|
|
|
// What makes a handled CUDA failure dangerous: the failed call leaves the error behind as the
|
|
// thread's last error, and nothing but cudaGetLastError() takes it away - later successful calls
|
|
// neither clear it nor return it. So a route that failed, was caught and was replaced by another
|
|
// one hands its error to the next cudaGetLastError() in the code that followed, which is the check
|
|
// after a kernel launch. That is how a broker logged a raw "out of memory" over an image whose
|
|
// buffers were all allocated. cuda_clear_error() is what a handled failure discharges it with.
|
|
TEST_CASE("CudaDevicePtr_FailedAllocationLeavesErrorUntilCleared", "[CUDAMemHelpers]") {
|
|
if (get_gpu_count() == 0)
|
|
SKIP("No CUDA GPU present");
|
|
|
|
cuda_clear_error();
|
|
REQUIRE_THROWS(CudaDevicePtr<uint8_t>(std::numeric_limits<size_t>::max(), CudaAlloc::Synchronous));
|
|
|
|
// Work that goes fine in between changes nothing: the error is still waiting.
|
|
CudaDevicePtr<uint8_t> fine(1 << 16);
|
|
REQUIRE(cudaStreamSynchronize(cuda_allocation_stream()) == cudaSuccess);
|
|
|
|
cuda_clear_error();
|
|
CHECK(cudaGetLastError() == cudaSuccess);
|
|
}
|
|
|
|
// The device decoder is allowed to fail - the host decompresses instead - but the host route runs
|
|
// kernels of its own and checks cudaGetLastError() after them. An image must come out of that
|
|
// whole, and with nothing left behind for the next one.
|
|
TEST_CASE("MXAnalysis_HandledDeviceDecodeFailureLeavesNoError", "[CUDAMemHelpers]") {
|
|
if (get_gpu_count() == 0)
|
|
SKIP("No CUDA GPU present");
|
|
|
|
constexpr size_t width = 1030, height = 1064;
|
|
|
|
DiffractionExperiment experiment(DetDECTRIS(width, height, "Test", {}));
|
|
experiment.ImagesPerTrigger(1).NumTriggers(1).PixelSigned(false)
|
|
.DetectorDistance_mm(75).BeamX_pxl(515).BeamY_pxl(532).IncidentEnergy_keV(12.4);
|
|
|
|
PixelMask pixel_mask(experiment);
|
|
AzimuthalIntegrationMapping mapping(experiment, pixel_mask);
|
|
IndexAndRefine index_and_refine(experiment, nullptr);
|
|
MXAnalysisWithoutFPGA analysis(experiment, mapping, pixel_mask, index_and_refine);
|
|
|
|
std::vector<uint16_t> pixels(width * height, 0);
|
|
for (size_t i = 0; i < pixels.size(); i += 997)
|
|
pixels[i] = 500;
|
|
JFJochBitShuffleCompressor compressor(CompressionAlgorithm::BSHUF_LZ4);
|
|
const std::vector<uint8_t> compressed = compressor.Compress(pixels);
|
|
REQUIRE(!compressed.empty());
|
|
|
|
DataMessage message{};
|
|
message.number = 0;
|
|
message.image = CompressedImage(compressed.data(), compressed.size(), width, height,
|
|
CompressedImageMode::Uint16, CompressionAlgorithm::BSHUF_LZ4);
|
|
|
|
SpotFindingSettings settings = DiffractionExperiment::DefaultDataProcessingSettings();
|
|
settings.enable = false;
|
|
|
|
// Stage the device route into failing, by leaving behind the error a request the card cannot
|
|
// serve leaves.
|
|
void *unservable = nullptr;
|
|
REQUIRE(cudaMalloc(&unservable, std::numeric_limits<size_t>::max()) != cudaSuccess);
|
|
|
|
AzimuthalIntegrationProfile profile(mapping);
|
|
REQUIRE_NOTHROW(analysis.Analyze(message, profile, settings));
|
|
CHECK(cudaGetLastError() == cudaSuccess);
|
|
}
|
|
|
|
#endif
|