Reset Cuda Context after exception

Question

I have a working app which uses Cuda / C++, but sometimes, because of memory leaks, throws exception. I need to be able to reset the GPU on live, my app is a server so it has to stay available.

I tried something like this, but it doesnt seems to work:

try
{
    // do process using GPU
}
catch (std::exception &e)
{
    // catching exception from cuda only

    cudaSetDevice(0);
    CUDA_RETURN_(cudaDeviceReset());
}

My idea is to reset the device each times I get an exception from the GPU, but I cannot manage to make it working. :( Btw, for some reasons, I cannot fix every problems of my Cuda code, I need a temporary solution. Thanks !

Capitalistic

Please clarify whether the memory leaks are on the GPU or not, and whether the exceptions are thrown by the CUDA runtime API or not. Also, why would memory leaks cause exceptions? At most you should encounter failures to allocate more memory. – Whittaker

Answer 1

The only method to restore proper device functionality after a non-recoverable ("sticky") CUDA error is to terminate the host process that initiated (i.e. issued the CUDA runtime API calls that led to) the error.

Therefore, for a single-process application, the only method is to terminate the application.

It should be possible to design a multi-process application, where the initial ("parent") process makes no usage of CUDA whatsoever, and spawns a child process that uses the GPU. When the child process encounters an unrecoverable CUDA error, it must terminate.

The parent process can, optionally, monitor the child process. If it determines that the child process has terminated, it can re-spawn the process and restore CUDA functional behavior.

Sticky vs. non-sticky errors are covered elsewhere, such as here.

An example of a proper multi-process app that uses e.g. fork() to spawn a child process that uses CUDA is available in the CUDA sample code simpleIPC. Here is a rough example assembled from the simpleIPC example (for linux):

$ cat t477.cu
/*
 * Copyright 1993-2015 NVIDIA Corporation.  All rights reserved.
 *
 * Please refer to the NVIDIA end user license agreement (EULA) associated
 * with this source code for terms and conditions that govern your use of
 * this software. Any use, reproduction, disclosure, or distribution of
 * this software and related documentation outside the terms of the EULA
 * is strictly prohibited.
 *
 */

// Includes
#include <stdio.h>
#include <assert.h>

// CUDA runtime includes
#include <cuda_runtime_api.h>

// CUDA utilities and system includes
#include <helper_cuda.h>

#define MAX_DEVICES          1
#define PROCESSES_PER_DEVICE 1
#define DATA_BUF_SIZE        4096

#ifdef __linux
#include <unistd.h>
#include <sched.h>
#include <sys/mman.h>
#include <sys/wait.h>
#include <linux/version.h>

typedef struct ipcDevices_st
{
    int count;
    int results[MAX_DEVICES];
} ipcDevices_t;


// CUDA Kernel
__global__ void simpleKernel(int *dst, int *src, int num)
{
    // Dummy kernel
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    dst[idx] = src[idx] / num;
}


void runTest(int index, ipcDevices_t* s_devices)
{
    if (s_devices->results[0] == 0){
        simpleKernel<<<1,1>>>(NULL, NULL, 1);  // make a fault
        cudaDeviceSynchronize();
        s_devices->results[0] = 1;}
    else {
        int *d, *s;
        int n = 1;
        cudaMalloc(&d, n*sizeof(int));
        cudaMalloc(&s, n*sizeof(int));
        simpleKernel<<<1,1>>>(d, s, n);
        cudaError_t err = cudaDeviceSynchronize();
        if (err != cudaSuccess)
          s_devices->results[0] = 0;
        else
          s_devices->results[0] = 2;}
    cudaDeviceReset();
}
#endif

int main(int argc, char **argv)
{

    ipcDevices_t *s_devices = (ipcDevices_t *) mmap(NULL, sizeof(*s_devices),
                                                    PROT_READ | PROT_WRITE, MAP_SHARED | MAP_ANONYMOUS, 0, 0);
    assert(MAP_FAILED != s_devices);

    // We can't initialize CUDA before fork() so we need to spawn a new process
    s_devices->count = 1;
    s_devices->results[0] = 0;

    printf("\nSpawning child process\n");
    int index = 0;

    pid_t pid = fork();

    printf("> Process %3d\n", pid);
    if (pid == 0) { // child process
    // launch our test
      runTest(index, s_devices);
    }
    // Cleanup and shutdown
    else { // parent process
            int status;
            waitpid(pid, &status, 0);
            if (s_devices->results[0] < 2) {
              printf("first process launch reported error: %d\n", s_devices->results[0]);
              printf("respawn\n");
              pid_t newpid = fork();
              if (newpid == 0) { // child process
                    // launch our test
                 runTest(index, s_devices);
                  }
    // Cleanup and shutdown
              else { // parent process
                int status;
                waitpid(newpid, &status, 0);
                if (s_devices->results[0] < 2)
                  printf("second process launch reported error: %d\n", s_devices->results[0]);
                else
                  printf("second process launch successful\n");
                }

            }

    }

    printf("\nShutting down...\n");

    exit(EXIT_SUCCESS);

}
$ nvcc -I/usr/local/cuda/samples/common/inc t477.cu -o t477
$ ./t477

Spawning child process
> Process 10841
> Process   0

Shutting down...
first process launch reported error: 1
respawn

Shutting down...
second process launch successful

Shutting down...
$

For windows, the only changes need should be to use a windows IPC mechanism for host interprocess communication.

Legere