Skip to content

Fall back to the cuDNN heuristics when benchmarking runs out of memory #3166

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

Merged
davisking merged 1 commit into from
Oct 3, 2026
+75 −12
Merged

Fall back to the cuDNN heuristics when benchmarking runs out of memory #3166

Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
87 changes: 75 additions & 12 deletions dlib/cuda/cudnn_dlibapi.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -786,6 +786,18 @@ namespace dlib
return best_alg;
}

// Whether cuDNN failed for lack of memory, which benchmarking the algorithms can run into
// when the rest of the program is using the GPU at the same time.
bool is_out_of_memory(cudnnStatus_t status)
{
#if CUDNN_MAJOR >= 9
return status == CUDNN_STATUS_INTERNAL_ERROR_HOST_ALLOCATION_FAILED ||
status == CUDNN_STATUS_INTERNAL_ERROR_DEVICE_ALLOCATION_FAILED;
#else
return status == CUDNN_STATUS_ALLOC_FAILED;
#endif
}

// The forward, backward data and backward filters algorithms chosen for each convolution
// configuration, with -1 for one not chosen yet. Once it has max_size configurations, it
// forgets the least recently used one for each new one.
Expand Down Expand Up @@ -841,7 +853,8 @@ namespace dlib
}

// Returns the algorithm that the cache has for the configuration, if cuDNN tells its workspace
// size, or else the one that find() returns, which then goes in the cache.
// size, or else the one that find() returns, which then goes in the cache - unless find()
// could not benchmark the algorithms and had to guess.
template <typename find_type, typename get_workspace_size_type>
int choose_algorithm (
const algorithm_cache::key_type& key,
Expand All @@ -867,9 +880,11 @@ namespace dlib
// fresh estimate, instead of using a cached value.
}
}
algorithms[which] = find();
bool benchmarked = true;
algorithms[which] = find(benchmarked);
workspace_size_in_bytes = get_workspace_size(algorithms[which]);
cache.insert(key, algorithms);
if (benchmarked)
cache.insert(key, algorithms);
return algorithms[which];
}

Expand All @@ -896,22 +911,38 @@ namespace dlib
if (forward_algo >= 0)
return;

const auto find = [&]() {
const auto find = [&](bool& benchmarked) {
cudnnConvolutionFwdAlgo_t forward_best_algo;
#if CUDNN_MAJOR >= 8
int num_possible_algorithms = 0;
CHECK_CUDNN(cudnnGetConvolutionForwardAlgorithmMaxCount(context(), &num_possible_algorithms));
std::vector<cudnnConvolutionFwdAlgoPerf_t> perf_results(num_possible_algorithms);
int num_algorithms = 0;
CHECK_CUDNN(cudnnFindConvolutionForwardAlgorithm(
cudnnStatus_t status = cudnnFindConvolutionForwardAlgorithm(
context(),
descriptor(data),
(const cudnnFilterDescriptor_t)filter_handle,
(const cudnnConvolutionDescriptor_t)conv_handle,
descriptor(output),
num_possible_algorithms,
&num_algorithms,
perf_results.data()));
perf_results.data());
// Benchmarking needs memory that the GPU may not have to spare at the moment, so
// fall back to the heuristics of cuDNN then.
if (is_out_of_memory(status))
{
benchmarked = false;
status = cudnnGetConvolutionForwardAlgorithm_v7(
context(),
descriptor(data),
(const cudnnFilterDescriptor_t)filter_handle,
(const cudnnConvolutionDescriptor_t)conv_handle,
descriptor(output),
num_possible_algorithms,
&num_algorithms,
perf_results.data());
}
CHECK_CUDNN(status);
perf_results.resize(num_algorithms);
forward_best_algo = pick_best_algorithm(perf_results);
#else
Expand Down Expand Up @@ -953,22 +984,38 @@ namespace dlib
if (backward_data_algo >= 0)
return;

const auto find = [&]() {
const auto find = [&](bool& benchmarked) {
cudnnConvolutionBwdDataAlgo_t backward_data_best_algo;
#if CUDNN_MAJOR >= 8
int num_possible_algorithms = 0;
CHECK_CUDNN(cudnnGetConvolutionBackwardFilterAlgorithmMaxCount(context(), &num_possible_algorithms));
std::vector<cudnnConvolutionBwdDataAlgoPerf_t> perf_results(num_possible_algorithms);
int num_algorithms = 0;
CHECK_CUDNN(cudnnFindConvolutionBackwardDataAlgorithm(
cudnnStatus_t status = cudnnFindConvolutionBackwardDataAlgorithm(
context(),
(const cudnnFilterDescriptor_t)filter_handle,
descriptor(gradient_input),
(const cudnnConvolutionDescriptor_t)conv_handle,
descriptor(data_gradient),
num_possible_algorithms,
&num_algorithms,
perf_results.data()));
perf_results.data());
// Benchmarking needs memory that the GPU may not have to spare at the moment, so
// fall back to the heuristics of cuDNN then.
if (is_out_of_memory(status))
{
benchmarked = false;
status = cudnnGetConvolutionBackwardDataAlgorithm_v7(
context(),
(const cudnnFilterDescriptor_t)filter_handle,
descriptor(gradient_input),
(const cudnnConvolutionDescriptor_t)conv_handle,
descriptor(data_gradient),
num_possible_algorithms,
&num_algorithms,
perf_results.data());
}
CHECK_CUDNN(status);
perf_results.resize(num_algorithms);
backward_data_best_algo = pick_best_algorithm(perf_results);
#else
Expand Down Expand Up @@ -1010,22 +1057,38 @@ namespace dlib
if (backward_filters_algo >= 0)
return;

const auto find = [&]() {
const auto find = [&](bool& benchmarked) {
cudnnConvolutionBwdFilterAlgo_t backward_filters_best_algo;
#if CUDNN_MAJOR >= 8
int num_possible_algorithms = 0;
CHECK_CUDNN(cudnnGetConvolutionBackwardFilterAlgorithmMaxCount(context(), &num_possible_algorithms));
std::vector<cudnnConvolutionBwdFilterAlgoPerf_t> perf_results(num_possible_algorithms);
int num_algorithms = 0;
CHECK_CUDNN(cudnnFindConvolutionBackwardFilterAlgorithm(
cudnnStatus_t status = cudnnFindConvolutionBackwardFilterAlgorithm(
context(),
descriptor(data),
descriptor(gradient_input),
(const cudnnConvolutionDescriptor_t)conv_handle,
(const cudnnFilterDescriptor_t)filter_handle,
num_possible_algorithms,
&num_algorithms,
perf_results.data()));
perf_results.data());
// Benchmarking needs memory that the GPU may not have to spare at the moment, so
// fall back to the heuristics of cuDNN then.
if (is_out_of_memory(status))
{
benchmarked = false;
status = cudnnGetConvolutionBackwardFilterAlgorithm_v7(
context(),
descriptor(data),
descriptor(gradient_input),
(const cudnnConvolutionDescriptor_t)conv_handle,
(const cudnnFilterDescriptor_t)filter_handle,
num_possible_algorithms,
&num_algorithms,
perf_results.data());
}
CHECK_CUDNN(status);
perf_results.resize(num_algorithms);
backward_filters_best_algo = pick_best_algorithm(perf_results);
#else
Expand Down
Loading