Skip to content

Fallback to different cuDNN algorithm when under memory pressure #2211

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
shelhamer merged 1 commit into from
Mar 26, 2015
Merged

Fallback to different cuDNN algorithm when under memory pressure #2211

Show file tree
Hide file tree
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
2 changes: 2 additions & 0 deletions src/caffe/layers/cudnn_conv_layer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,6 +24,8 @@ void CuDNNConvolutionLayer<Dtype>::LayerSetUp(
// Initialize CUDA streams and cuDNN.
stream_ = new cudaStream_t[this->group_ * CUDNN_STREAMS_PER_GROUP];
handle_ = new cudnnHandle_t[this->group_ * CUDNN_STREAMS_PER_GROUP];
workspaceSizeInBytes = 0;
workspace = NULL;

for (int g = 0; g < this->group_ * CUDNN_STREAMS_PER_GROUP; g++) {
CUDA_CHECK(cudaStreamCreate(&stream_[g]));
Expand Down
19 changes: 15 additions & 4 deletions src/caffe/layers/cudnn_conv_layer.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,6 +19,11 @@ void CuDNNConvolutionLayer<Dtype>::Forward_gpu(
Dtype* top_data = top[i]->mutable_gpu_data();
const Dtype* weight = this->blobs_[0]->gpu_data();

size_t workspace_limit_bytes = this->kernel_h_ *

@shelhamer shelhamer Mar 26, 2015

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Will this always rule out the cuDNN GEMM convolution? At least in the Caffe GEMM convolution the workspace is the kernel dimensions (kernel_h * kernel_w * channels) * output dimensions (height_out_* width_out_) as in https://github.com/BVLC/caffe/blob/master/src/caffe/layers/base_conv_layer.cpp#L147, although in the cuDNN implementation I suppose the workspace could be just the input data so the +1 allows it here.

@nsubtil nsubtil Mar 26, 2015

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, my understanding is that the workspace size depends only on the input data size.

The intent is that GEMM convolution will still be chosen if possible. However, in practice, we expect that there won't be enough memory available in many use cases.

this->kernel_w_ *
this->channels_ *
sizeof(int) + 1;

@shelhamer shelhamer Mar 26, 2015

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

sizeof(Dtype)?

@nsubtil nsubtil Mar 26, 2015

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actually, sizeof(int) is correct here. The workspace is always an integer buffer.

@shelhamer shelhamer Mar 26, 2015

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It seems odd to me that the workspace can be the same size whether the input and filters are float or double, but if that's how it is in cuDNN then so be it.


// Forward through cuDNN in parallel over groups.
for (int g = 0; g < this->group_; g++) {
cudnnConvolutionFwdAlgo_t algo;
Expand All @@ -32,8 +37,8 @@ void CuDNNConvolutionLayer<Dtype>::Forward_gpu(
filter_desc_,
conv_descs_[i],
top_descs_[i],
CUDNN_CONVOLUTION_FWD_PREFER_FASTEST,
0, // memoryLimitInBytes,
CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT,
workspace_limit_bytes, // memoryLimitInBytes,
&algo));

// get minimum size of the workspace needed for the desired algorithm
Expand All @@ -45,13 +50,19 @@ void CuDNNConvolutionLayer<Dtype>::Forward_gpu(
conv_descs_[i],
top_descs_[i],
algo,
&workspaceSizeInBytes));
&workspaceSizeInBytes_temp));

if (workspaceSizeInBytes_temp > workspaceSizeInBytes) {
workspaceSizeInBytes = workspaceSizeInBytes_temp;
// free the existing workspace and allocate a new (larger) one
cudaFree(this->workspace);
cudaMalloc(&(this->workspace), workspaceSizeInBytes);
cudaError_t err = cudaMalloc(&(this->workspace), workspaceSizeInBytes);
if (err != cudaSuccess) {
// force zero memory path
algo = CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM;
workspace = NULL;
workspaceSizeInBytes = 0;
}
}

// Filters.
Expand Down