Skip to content

Commit 2d25773

Browse files
davidrohrdavidrohr
committed
GPU: In case of segfault on device, print stored error codes before throwing
1 parent d1d0995 commit 2d25773

15 files changed

Lines changed: 112 additions & 88 deletions

‎GPU/GPUTracking/Base/GPUReconstruction.cxx‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1000,11 +1000,11 @@ void GPUReconstruction::PrepareEvent() // TODO: Clean this up, this should not b
10001000
AllocateRegisteredMemory(nullptr);
10011001
}
10021002

1003-
int GPUReconstruction::CheckErrorCodes(bool cpuOnly)
1003+
int GPUReconstruction::CheckErrorCodes(bool cpuOnly, bool forceShowErrors)
10041004
{
10051005
int retVal = 0;
10061006
for (unsigned int i = 0; i < mChains.size(); i++) {
1007-
if (mChains[i]->CheckErrorCodes(cpuOnly)) {
1007+
if (mChains[i]->CheckErrorCodes(cpuOnly, forceShowErrors)) {
10081008
retVal++;
10091009
}
10101010
}

‎GPU/GPUTracking/Base/GPUReconstruction.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -189,7 +189,7 @@ class GPUReconstruction
189189
virtual void* getGPUPointer(void* ptr) { return ptr; }
190190
virtual void startGPUProfiling() {}
191191
virtual void endGPUProfiling() {}
192-
int CheckErrorCodes(bool cpuOnly = false);
192+
int CheckErrorCodes(bool cpuOnly = false, bool forceShowErrors = false);
193193
void RunPipelineWorker();
194194
void TerminatePipelineWorker();
195195

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu‎

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -77,6 +77,28 @@ GPUReconstructionCUDABackend::~GPUReconstructionCUDABackend()
7777
}
7878
}
7979

80+
int GPUReconstructionCUDABackend::GPUFailedMsgAI(const long long int error, const char* file, int line)
81+
{
82+
// Check for CUDA Error and in the case of an error display the corresponding error string
83+
if (error == cudaSuccess) {
84+
return (0);
85+
}
86+
GPUError("CUDA Error: %lld / %s (%s:%d)", error, cudaGetErrorString((cudaError_t)error), file, line);
87+
return 1;
88+
}
89+
90+
void GPUReconstructionCUDABackend::GPUFailedMsgA(const long long int error, const char* file, int line)
91+
{
92+
if (GPUFailedMsgAI(error, file, line)) {
93+
static bool runningCallbacks = false;
94+
if (IsInitialized() && runningCallbacks == false) {
95+
runningCallbacks = true;
96+
CheckErrorCodes(false, true);
97+
}
98+
throw std::runtime_error("CUDA Failure");
99+
}
100+
}
101+
80102
GPUReconstructionCUDA::GPUReconstructionCUDA(const GPUSettingsDeviceBackend& cfg) : GPUReconstructionKernels(cfg)
81103
{
82104
mDeviceBackendSettings.deviceType = DeviceType::CUDA;

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.h‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -33,6 +33,8 @@ class GPUReconstructionCUDABackend : public GPUReconstructionDeviceBase
3333
{
3434
public:
3535
~GPUReconstructionCUDABackend() override;
36+
int GPUFailedMsgAI(const long long int error, const char* file, int line);
37+
void GPUFailedMsgA(const long long int error, const char* file, int line);
3638

3739
protected:
3840
GPUReconstructionCUDABackend(const GPUSettingsDeviceBackend& cfg);

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAInternals.h‎

Lines changed: 0 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -40,23 +40,6 @@ struct GPUReconstructionCUDAInternals {
4040
#define GPUFailedMsg(x) GPUFailedMsgA(x, __FILE__, __LINE__)
4141
#define GPUFailedMsgI(x) GPUFailedMsgAI(x, __FILE__, __LINE__)
4242

43-
static int GPUFailedMsgAI(const long long int error, const char* file, int line)
44-
{
45-
// Check for CUDA Error and in the case of an error display the corresponding error string
46-
if (error == cudaSuccess) {
47-
return (0);
48-
}
49-
GPUError("CUDA Error: %lld / %s (%s:%d)", error, cudaGetErrorString((cudaError_t)error), file, line);
50-
return 1;
51-
}
52-
53-
static void GPUFailedMsgA(const long long int error, const char* file, int line)
54-
{
55-
if (GPUFailedMsgAI(error, file, line)) {
56-
throw std::runtime_error("CUDA Failure");
57-
}
58-
}
59-
6043
static_assert(std::is_convertible<cudaEvent_t, void*>::value, "CUDA event type incompatible to deviceEvent");
6144

6245
} // namespace gpu

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAKernels.cu‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -29,11 +29,11 @@ texture<calink, cudaTextureType1D, cudaReadModeElementType> gAliTexRefu;
2929
class GPUDebugTiming
3030
{
3131
public:
32-
GPUDebugTiming(bool d, void** t, cudaStream_t* s, GPUReconstruction::krnlSetup& x, GPUReconstructionCUDABackend* r = nullptr) : mDeviceTimers(t), mStreams(s), mXYZ(x), mRec(r), mDo(d)
32+
GPUDebugTiming(bool d, void** t, cudaStream_t* s, GPUReconstruction::krnlSetup& x, GPUReconstructionCUDABackend* r) : mDeviceTimers(t), mStreams(s), mXYZ(x), mRec(r), mDo(d)
3333
{
3434
if (mDo) {
3535
if (mDeviceTimers) {
36-
GPUFailedMsg(cudaEventRecord((cudaEvent_t)mDeviceTimers[0], mStreams[mXYZ.x.stream]));
36+
mRec->GPUFailedMsg(cudaEventRecord((cudaEvent_t)mDeviceTimers[0], mStreams[mXYZ.x.stream]));
3737
} else {
3838
mTimer.ResetStart();
3939
}
@@ -43,13 +43,13 @@ class GPUDebugTiming
4343
{
4444
if (mDo) {
4545
if (mDeviceTimers) {
46-
GPUFailedMsg(cudaEventRecord((cudaEvent_t)mDeviceTimers[1], mStreams[mXYZ.x.stream]));
47-
GPUFailedMsg(cudaEventSynchronize((cudaEvent_t)mDeviceTimers[1]));
46+
mRec->GPUFailedMsg(cudaEventRecord((cudaEvent_t)mDeviceTimers[1], mStreams[mXYZ.x.stream]));
47+
mRec->GPUFailedMsg(cudaEventSynchronize((cudaEvent_t)mDeviceTimers[1]));
4848
float v;
49-
GPUFailedMsg(cudaEventElapsedTime(&v, (cudaEvent_t)mDeviceTimers[0], (cudaEvent_t)mDeviceTimers[1]));
49+
mRec->GPUFailedMsg(cudaEventElapsedTime(&v, (cudaEvent_t)mDeviceTimers[0], (cudaEvent_t)mDeviceTimers[1]));
5050
mXYZ.t = v * 1.e-3;
5151
} else {
52-
GPUFailedMsg(cudaStreamSynchronize(mStreams[mXYZ.x.stream]));
52+
mRec->GPUFailedMsg(cudaStreamSynchronize(mStreams[mXYZ.x.stream]));
5353
mXYZ.t = mTimer.GetCurrentElapsedTime();
5454
}
5555
}
@@ -117,7 +117,7 @@ static void getArgPtrs(const void** pArgs, const T& arg, const Args&... args)
117117
template <class T, int I, typename... Args>
118118
void GPUReconstructionCUDABackend::runKernelBackendInternal(krnlSetup& _xyz, const Args&... args)
119119
{
120-
GPUDebugTiming timer(mProcessingSettings.deviceTimers && mProcessingSettings.debugLevel > 0, (void**)mDebugEvents, mInternals->Streams, _xyz);
120+
GPUDebugTiming timer(mProcessingSettings.deviceTimers && mProcessingSettings.debugLevel > 0, (void**)mDebugEvents, mInternals->Streams, _xyz, this);
121121
if (mProcessingSettings.rtc.enable) {
122122
auto& x = _xyz.x;
123123
auto& y = _xyz.y;

‎GPU/GPUTracking/Base/hip/GPUReconstructionHIP.h‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -33,6 +33,8 @@ class GPUReconstructionHIPBackend : public GPUReconstructionDeviceBase
3333
{
3434
public:
3535
~GPUReconstructionHIPBackend() override;
36+
int GPUFailedMsgAI(const long long int error, const char* file, int line);
37+
void GPUFailedMsgA(const long long int error, const char* file, int line);
3638

3739
protected:
3840
GPUReconstructionHIPBackend(const GPUSettingsDeviceBackend& cfg);

‎GPU/GPUTracking/Base/hip/GPUReconstructionHIP.hip.cxx‎

Lines changed: 28 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -79,11 +79,11 @@ class TimeFrameGPU : public TimeFrame
7979
class GPUDebugTiming
8080
{
8181
public:
82-
GPUDebugTiming(bool d, void** t, hipStream_t* s, GPUReconstruction::krnlSetup& x, GPUReconstructionHIPBackend* r = nullptr) : mDeviceTimers(t), mStreams(s), mXYZ(x), mRec(r), mDo(d)
82+
GPUDebugTiming(bool d, void** t, hipStream_t* s, GPUReconstruction::krnlSetup& x, GPUReconstructionHIPBackend* r) : mDeviceTimers(t), mStreams(s), mXYZ(x), mRec(r), mDo(d)
8383
{
8484
if (mDo) {
8585
if (mDeviceTimers) {
86-
GPUFailedMsg(hipEventRecord((hipEvent_t)mDeviceTimers[0], mStreams[mXYZ.x.stream]));
86+
mRec->GPUFailedMsg(hipEventRecord((hipEvent_t)mDeviceTimers[0], mStreams[mXYZ.x.stream]));
8787
} else {
8888
mTimer.ResetStart();
8989
}
@@ -93,13 +93,13 @@ class GPUDebugTiming
9393
{
9494
if (mDo) {
9595
if (mDeviceTimers) {
96-
GPUFailedMsg(hipEventRecord((hipEvent_t)mDeviceTimers[1], mStreams[mXYZ.x.stream]));
97-
GPUFailedMsg(hipEventSynchronize((hipEvent_t)mDeviceTimers[1]));
96+
mRec->GPUFailedMsg(hipEventRecord((hipEvent_t)mDeviceTimers[1], mStreams[mXYZ.x.stream]));
97+
mRec->GPUFailedMsg(hipEventSynchronize((hipEvent_t)mDeviceTimers[1]));
9898
float v;
99-
GPUFailedMsg(hipEventElapsedTime(&v, (hipEvent_t)mDeviceTimers[0], (hipEvent_t)mDeviceTimers[1]));
99+
mRec->GPUFailedMsg(hipEventElapsedTime(&v, (hipEvent_t)mDeviceTimers[0], (hipEvent_t)mDeviceTimers[1]));
100100
mXYZ.t = v * 1.e-3;
101101
} else {
102-
GPUFailedMsg(hipStreamSynchronize(mStreams[mXYZ.x.stream]));
102+
mRec->GPUFailedMsg(hipStreamSynchronize(mStreams[mXYZ.x.stream]));
103103
mXYZ.t = mTimer.GetCurrentElapsedTime();
104104
}
105105
}
@@ -211,6 +211,28 @@ GPUReconstructionHIPBackend::~GPUReconstructionHIPBackend()
211211
}
212212
}
213213

214+
int GPUReconstructionHIPBackend::GPUFailedMsgAI(const long long int error, const char* file, int line)
215+
{
216+
// Check for HIP Error and in the case of an error display the corresponding error string
217+
if (error == hipSuccess) {
218+
return (0);
219+
}
220+
GPUError("HIP Error: %lld / %s (%s:%d)", error, hipGetErrorString((hipError_t)error), file, line);
221+
return 1;
222+
}
223+
224+
void GPUReconstructionHIPBackend::GPUFailedMsgA(const long long int error, const char* file, int line)
225+
{
226+
if (GPUFailedMsgAI(error, file, line)) {
227+
static bool runningCallbacks = false;
228+
if (IsInitialized() && runningCallbacks == false) {
229+
runningCallbacks = true;
230+
CheckErrorCodes(false, true);
231+
}
232+
throw std::runtime_error("HIP Failure");
233+
}
234+
}
235+
214236
GPUReconstruction* GPUReconstruction_Create_HIP(const GPUSettingsDeviceBackend& cfg) { return new GPUReconstructionHIP(cfg); }
215237

216238
void GPUReconstructionHIPBackend::GetITSTraits(std::unique_ptr<o2::its::TrackerTraits>* trackerTraits, std::unique_ptr<o2::its::VertexerTraits>* vertexerTraits)

‎GPU/GPUTracking/Base/hip/GPUReconstructionHIPInternals.h‎

Lines changed: 0 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -30,23 +30,6 @@ struct GPUReconstructionHIPInternals {
3030
#define GPUFailedMsg(x) GPUFailedMsgA(x, __FILE__, __LINE__)
3131
#define GPUFailedMsgI(x) GPUFailedMsgAI(x, __FILE__, __LINE__)
3232

33-
static int GPUFailedMsgAI(const long long int error, const char* file, int line)
34-
{
35-
// Check for HIP Error and in the case of an error display the corresponding error string
36-
if (error == hipSuccess) {
37-
return (0);
38-
}
39-
GPUError("HIP Error: %lld / %s (%s:%d)", error, hipGetErrorString((hipError_t)error), file, line);
40-
return 1;
41-
}
42-
43-
static void GPUFailedMsgA(const long long int error, const char* file, int line)
44-
{
45-
if (GPUFailedMsgAI(error, file, line)) {
46-
throw std::runtime_error("HIP Failure");
47-
}
48-
}
49-
5033
static_assert(std::is_convertible<hipEvent_t, void*>::value, "HIP event type incompatible to deviceEvent");
5134

5235
} // namespace gpu

‎GPU/GPUTracking/Base/opencl-common/GPUReconstructionOCL.cxx‎

Lines changed: 25 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -53,6 +53,28 @@ GPUReconstructionOCL::~GPUReconstructionOCL()
5353
}
5454
}
5555

56+
int GPUReconstructionOCL::GPUFailedMsgAI(const long int error, const char* file, int line)
57+
{
58+
// Check for OPENCL Error and in the case of an error display the corresponding error string
59+
if (error == CL_SUCCESS) {
60+
return (0);
61+
}
62+
GPUError("OCL Error: %ld / %s (%s:%d)", error, opencl_error_string(error), file, line);
63+
return 1;
64+
}
65+
66+
void GPUReconstructionOCL::GPUFailedMsgA(const long int error, const char* file, int line)
67+
{
68+
if (GPUFailedMsgAI(error, file, line)) {
69+
static bool runningCallbacks = false;
70+
if (IsInitialized() && runningCallbacks == false) {
71+
runningCallbacks = true;
72+
CheckErrorCodes(false, true);
73+
}
74+
throw std::runtime_error("OpenCL Failure");
75+
}
76+
}
77+
5678
void GPUReconstructionOCL::UpdateSettings()
5779
{
5880
GPUCA_GPUReconstructionUpdateDefailts();
@@ -296,8 +318,9 @@ int GPUReconstructionOCL::InitDevice_Runtime()
296318
if (GPUFailedMsgI(ocl_error)) {
297319
quit("Error creating kernel");
298320
}
299-
OCLsetKernelParameters(kernel, mInternals->mem_gpu, mInternals->mem_constant, mInternals->mem_host);
300-
if (GPUFailedMsgI(clExecuteKernelA(mInternals->command_queue[0], kernel, 16, 16, nullptr)) ||
321+
322+
if (GPUFailedMsgI(OCLsetKernelParameters(kernel, mInternals->mem_gpu, mInternals->mem_constant, mInternals->mem_host)) ||
323+
GPUFailedMsgI(clExecuteKernelA(mInternals->command_queue[0], kernel, 16, 16, nullptr)) ||
301324
GPUFailedMsgI(clFinish(mInternals->command_queue[0])) ||
302325
GPUFailedMsgI(clReleaseKernel(kernel)) ||
303326
GPUFailedMsgI(clReleaseProgram(program))) {

0 commit comments

Comments
 (0)