Skip to content

Improve GPU error handling #816

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
TysonRayJones wants to merge 4 commits into
base: devel
Choose a base branch
from
Open

Improve GPU error handling #816

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
c36af60
add kernel launch error handling
TysonRayJones Aug 17, 2026
3454426
corrected cuQuantum error macro
TysonRayJones Aug 17, 2026
4c4d700
tidy
TysonRayJones Aug 17, 2026
e7cfbf9
catch Thrust errors
TysonRayJones Aug 17, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
35 changes: 30 additions & 5 deletions quest/src/core/errors.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -694,11 +694,13 @@ void assert_gpuHasBeenBound(bool isBound) {
* CUDA ERRORS
*/

void error_cudaCallFailed(const char* msg, const char* func, const char* caller, const char* file, int line) {
void internal_cudaLibCallFailed(const char* libname, const char* msg, const char* func, const char* caller, const char* file, int line) {

// using operator overloads to cast const char[] literals to std::string, to concat with const char*.
string err = "";
err += "A CUDA (or cuQuantum) API function (\"";
err += "A ";
err += libname;
err += " API function (\"";
err += func;
err += "\", called by \"";
err += caller;
Expand All @@ -712,21 +714,39 @@ void error_cudaCallFailed(const char* msg, const char* func, const char* caller,
raiseInternalError(err);
}

void error_cudaCallFailed(const char* msg, const char* func, const char* caller, const char* file, int line) {

internal_cudaLibCallFailed("CUDA", msg, func, caller, file, line);
}

void error_cudaEncounteredIrrecoverableError() {

raiseInternalError("The CUDA API encountered an irrecoverable \"sticky\" error which was attemptedly cleared as if it were non-sticky.");
}

void error_cudaKernelLaunchFailed(const char* caller, const char* cudaErrMsg) {

string err = "";
err += "A CUDA kernel invoked within '";
err += caller;
err += "' failed to launch - or a prior kernel called from elsewhere asynchronously failed -";
err += " with CUDA error message: \"";
err += cudaErrMsg;
err += "\". ";
raiseInternalError(err);
}

// Looking for assert_lastKernelLaunchSucceeded(const char*)? It's in gpu_config :^)



/*
* THRUST ERRORS
*/

void error_thrustCallFailed(const char* msg, const char* func, const char* caller, const char* file, int line) {

void error_thrustTempGpuAllocFailed() {

raiseInternalError("Thrust failed to allocate temporary GPU memory.");
internal_cudaLibCallFailed("Thrust", msg, func, caller, file, line);
}


Expand All @@ -735,6 +755,11 @@ void error_thrustTempGpuAllocFailed() {
* CUQUANTUM ERRORS
*/

void error_cuQuantumCallFailed(const char* msg, const char* func, const char* caller, const char* file, int line) {

internal_cudaLibCallFailed("cuQuantum (specifically cuStateVec)", msg, func, caller, file, line);
}

void error_cuQuantumInitOrFinalizedButNotCompiled() {

raiseInternalError("Attempted to initialise or finalise cuQuantum, but cuQuantum was not compiled.");
Expand Down
8 changes: 7 additions & 1 deletion quest/src/core/errors.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -281,20 +281,26 @@ void error_cudaCallFailed(const char* msg, const char* func, const char* caller,

void error_cudaEncounteredIrrecoverableError();

void error_cudaKernelLaunchFailed(const char* caller, const char* cudaErrMsg);

// Looking for assert_lastKernelLaunchSucceeded(const char*)? It's in gpu_config :^)



/*
* THRUST ERRORS
*/

void error_thrustTempGpuAllocFailed();
void error_thrustCallFailed(const char* msg, const char* func, const char* caller, const char* file, int line);



/*
* CUQUANTUM ERRORS
*/

void error_cuQuantumCallFailed(const char* msg, const char* func, const char* caller, const char* file, int line);

void error_cuQuantumInitOrFinalizedButNotCompiled();

void error_cuQuantumTempCpuAllocFailed();
Expand Down
13 changes: 12 additions & 1 deletion quest/src/gpu/gpu_config.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -91,6 +91,17 @@ void clearPossibleCudaError() {
error_cudaEncounteredIrrecoverableError();
}

void assert_lastKernelLaunchSucceeded(const char* caller) {

// note that we are only checking the kernel LAUNCH succeeded; it remains
// possible for the kernel to subsequently fail, which would only be detected
// at a subsequent cudaGetLastError(), or a cudaDeviceSynchronize().

cudaError_t status = cudaGetLastError();
if (status != cudaSuccess)
error_cudaKernelLaunchFailed(caller, cudaGetErrorString(status));
}

#endif


Expand Down Expand Up @@ -360,7 +371,7 @@ int gpu_getMaxNumThreadsPerBlock() {
#if QUEST_COMPILE_CUDA

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, getBoundGpuId());
CUDA_CHECK( cudaGetDeviceProperties(&prop, getBoundGpuId()) );
return prop.maxThreadsPerBlock; // HIP compatible

#else
Expand Down
2 changes: 2 additions & 0 deletions quest/src/gpu/gpu_config.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,6 +40,8 @@ constexpr int gpu_HIP_WARP_SIZE = 64;

void assertCudaCallSucceeded(int code, const char* call, const char* caller, const char* file, int line);

void assert_lastKernelLaunchSucceeded(const char* caller);

#endif


Expand Down
42 changes: 29 additions & 13 deletions quest/src/gpu/gpu_cuquantum.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,6 +44,7 @@

#include "quest/include/precision.h"

#include "quest/src/core/errors.hpp"
#include "quest/src/core/lists.hpp"
#include "quest/src/core/utilities.hpp"
#include "quest/src/gpu/gpu_config.hpp"
Expand Down Expand Up @@ -77,6 +78,21 @@ using std::vector;



/*
* CUSTATEVEC ERROR HANDLING
*/

inline void assertCuStateVecCallSucceeded(custatevecStatus_t status, const char* call, const char* caller, const char* file, int line) {

if (status != CUSTATEVEC_STATUS_SUCCESS)
error_cuQuantumCallFailed(custatevecGetErrorString(status), call, caller, file, line);
}

#define CUSV_CHECK(cmd) \
assertCuStateVecCallSucceeded((cmd), #cmd, __func__, __FILE__, __LINE__)



/*
* ENVIRONMENT MANAGEMENT
*/
Expand Down Expand Up @@ -143,7 +159,7 @@ void gpu_initCuQuantum() {
// prior validation prevent it (disabled by an environment variable)

// create new stream and cuQuantum handle, binding to global config
CUDA_CHECK( custatevecCreate(&config.handle) );
CUSV_CHECK( custatevecCreate(&config.handle) );
CUDA_CHECK( cudaStreamCreate(&config.stream) );

// get and configure existing memory pool (for later automatic alloc/dealloc of gate matrices)
Expand All @@ -157,15 +173,15 @@ void gpu_initCuQuantum() {
strcpy(config.memhandler.name, "mempool");

// bind memory handler and stream to cuQuantum handle
CUDA_CHECK( custatevecSetDeviceMemHandler(config.handle, &config.memhandler) );
CUDA_CHECK( custatevecSetStream(config.handle, config.stream) );
CUSV_CHECK( custatevecSetDeviceMemHandler(config.handle, &config.memhandler) );
CUSV_CHECK( custatevecSetStream(config.handle, config.stream) );
}


void gpu_finalizeCuQuantum() {

CUDA_CHECK( cudaStreamDestroy(config.stream) );
CUDA_CHECK( custatevecDestroy(config.handle) );
CUSV_CHECK( custatevecDestroy(config.handle) );
}


Expand All @@ -181,7 +197,7 @@ void cuquantum_statevec_anyCtrlSwap_subA(Qureg qureg, ConstList64 ctrls, ConstLi
int2 targPairs[] = {{targ1, targ2}};;
int numTargPairs = 1;

CUDA_CHECK( custatevecSwapIndexBits(
CUSV_CHECK( custatevecSwapIndexBits(
config.handle,
getGpuQcompPtr(qureg.gpuAmps), CUQUANTUM_QCOMP, qureg.logNumAmpsPerNode,
targPairs, numTargPairs,
Expand Down Expand Up @@ -209,7 +225,7 @@ void cuquantum_statevec_anyCtrlAnyTargDenseMatrix_subA(Qureg qureg, ConstList64
void* work = nullptr;
size_t workSize = 0;

CUDA_CHECK( custatevecApplyMatrix(
CUSV_CHECK( custatevecApplyMatrix(
config.handle,
getGpuQcompPtr(qureg.gpuAmps), CUQUANTUM_QCOMP, qureg.logNumAmpsPerNode,
flatMatrElems, CUQUANTUM_QCOMP, CUSTATEVEC_MATRIX_LAYOUT_ROW, applyAdj,
Expand Down Expand Up @@ -238,7 +254,7 @@ void cuquantum_statevec_anyCtrlAnyTargDiagMatr_sub(Qureg qureg, ConstList64 ctrl
void* work = nullptr;
size_t workSize = 0;

CUDA_CHECK( custatevecApplyGeneralizedPermutationMatrix(
CUSV_CHECK( custatevecApplyGeneralizedPermutationMatrix(
config.handle,
getGpuQcompPtr(qureg.gpuAmps), CUQUANTUM_QCOMP, qureg.logNumAmpsPerNode,
perm, flatMatrElems, CUQUANTUM_QCOMP, adj,
Expand Down Expand Up @@ -335,7 +351,7 @@ qreal cuquantum_statevec_calcTotalProb_sub(Qureg qureg) {
int qubit = qureg.logNumAmpsPerNode - 1;
int numQubits = 1;

CUDA_CHECK( custatevecAbs2SumOnZBasis(
CUSV_CHECK( custatevecAbs2SumOnZBasis(
config.handle,
getGpuQcompPtr(qureg.gpuAmps), CUQUANTUM_QCOMP, qureg.logNumAmpsPerNode,
&prob0, &prob1, &qubit, numQubits ) );
Expand All @@ -350,7 +366,7 @@ qreal cuquantum_statevec_calcProbOfMultiQubitOutcome_sub(Qureg qureg, ConstList6
// cuQuantum probabilities are always double
double prob;

CUDA_CHECK( custatevecAbs2SumArray(
CUSV_CHECK( custatevecAbs2SumArray(
config.handle,
getGpuQcompPtr(qureg.gpuAmps), CUQUANTUM_QCOMP, qureg.logNumAmpsPerNode,
&prob, nullptr, 0, outcomes.data(), qubits.data(), qubits.size()) );
Expand All @@ -371,7 +387,7 @@ void cuquantum_statevec_calcProbsOfAllMultiQubitOutcomes_sub(qreal* outProbs, Qu
double* outPtr = tmpProbs.data();
#endif

CUDA_CHECK( custatevecAbs2SumArray(
CUSV_CHECK( custatevecAbs2SumArray(
config.handle,
getGpuQcompPtr(qureg.gpuAmps), CUQUANTUM_QCOMP, qureg.logNumAmpsPerNode,
outPtr, qubits.data(), qubits.size(), nullptr, nullptr, 0) );
Expand Down Expand Up @@ -413,7 +429,7 @@ qreal cuquantum_statevec_calcExpecPauliStr_subA(Qureg qureg, ConstList64 x, Cons
// cuStateVec output is always double
double value = 0;

CUDA_CHECK( custatevecComputeExpectationsOnPauliBasis(
CUSV_CHECK( custatevecComputeExpectationsOnPauliBasis(
config.handle,
getGpuQcompPtr(qureg.gpuAmps), CUQUANTUM_QCOMP, qureg.logNumAmpsPerNode,
&value, termPaulis, numTerms, termTargets, numPaulisPerTerm) );
Expand All @@ -437,11 +453,11 @@ qreal cuquantum_statevec_calcExpecAnyTargZ_sub(Qureg qureg, ConstList64 targs) {

void cuquantum_statevec_multiQubitProjector_sub(Qureg qureg, ConstList64 qubits, ConstList64 outcomes, qreal prob) {

CUDA_CHECK( custatevecCollapseByBitString(
CUSV_CHECK( custatevecCollapseByBitString(
config.handle,
getGpuQcompPtr(qureg.gpuAmps), CUQUANTUM_QCOMP, qureg.logNumAmpsPerNode,
outcomes.data(), qubits.data(), qubits.size(), prob) );
}


#endif // GPU_CUQUANTUM_HPP
#endif // GPU_CUQUANTUM_HPP
Loading
Loading