| 2068 | |
| 2069 | template <class T> |
| 2070 | port::Status CudnnSupport::DoRnnBackwardImpl( |
| 2071 | Stream* stream, const CudnnRnnDescriptor& rnn_desc, |
| 2072 | const CudnnRnnSequenceTensorDescriptor& input_desc, |
| 2073 | const DeviceMemory<T>& input_data, |
| 2074 | const CudnnRnnStateTensorDescriptor& input_h_desc, |
| 2075 | const DeviceMemory<T>& input_h_data, |
| 2076 | const CudnnRnnStateTensorDescriptor& input_c_desc, |
| 2077 | const DeviceMemory<T>& input_c_data, const DeviceMemory<T>& params, |
| 2078 | const CudnnRnnSequenceTensorDescriptor& output_desc, |
| 2079 | const DeviceMemory<T>& output_data, |
| 2080 | const CudnnRnnStateTensorDescriptor& output_h_desc, |
| 2081 | const DeviceMemory<T>& output_h_data, |
| 2082 | const CudnnRnnStateTensorDescriptor& output_c_desc, |
| 2083 | const DeviceMemory<T>& output_c_data, |
| 2084 | const DeviceMemory<T>& output_backprop_data, |
| 2085 | const DeviceMemory<T>& output_h_backprop_data, |
| 2086 | const DeviceMemory<T>& output_c_backprop_data, |
| 2087 | DeviceMemory<T>* input_backprop_data, |
| 2088 | DeviceMemory<T>* input_h_backprop_data, |
| 2089 | DeviceMemory<T>* input_c_backprop_data, |
| 2090 | DeviceMemory<T>* params_backprop_data, |
| 2091 | DeviceMemory<uint8>* reserve_space_data, |
| 2092 | ScratchAllocator* workspace_allocator, |
| 2093 | dnn::ProfileResult* output_profile_result) { |
| 2094 | SE_ASSIGN_OR_RETURN( |
| 2095 | RnnModelDims model_dims, |
| 2096 | ExtractAndCheckRnnForward(rnn_desc, input_desc, input_data, input_h_desc, |
| 2097 | input_h_data, input_c_desc, input_c_data, |
| 2098 | params, output_desc, output_data, output_h_desc, |
| 2099 | output_h_data, output_c_desc, output_c_data)); |
| 2100 | |
| 2101 | auto cudnn = cudnn_->GetHandle(parent_, stream); |
| 2102 | |
| 2103 | SE_RETURN_IF_ERROR(CheckRNNParameterSize(cudnn, rnn_desc, input_desc)); |
| 2104 | SE_ASSIGN_OR_RETURN(DeviceMemory<uint8> workspace, |
| 2105 | CreateRnnWorkspace(stream, cudnn, rnn_desc, input_desc, |
| 2106 | workspace_allocator)); |
| 2107 | |
| 2108 | std::unique_ptr<GpuTimer, GpuTimerDeleter> timer; |
| 2109 | const bool is_profiling = output_profile_result != nullptr; |
| 2110 | if (is_profiling) { |
| 2111 | timer.reset(new GpuTimer(parent_)); |
| 2112 | // The start and stop of the timer should be as close to the Cudnn call as |
| 2113 | // possible. It is still possible for other threads to issue workload on |
| 2114 | // to this stream. So it could take multiple profiling measurements. |
| 2115 | if (!timer->Init() || !timer->Start(AsGpuStream(stream))) { |
| 2116 | return port::Status(port::error::INTERNAL, "Failed to start timer"); |
| 2117 | } |
| 2118 | } |
| 2119 | |
| 2120 | if (input_desc.is_var_seq_lengths()) { |
| 2121 | #if CUDNN_VERSION >= 7201 |
| 2122 | RETURN_IF_CUDNN_ERROR(cudnnRNNBackwardDataEx( |
| 2123 | /*handle=*/cudnn.handle(), /*rnnDesc=*/rnn_desc.handle(), |
| 2124 | /*yDesc=*/output_desc.data_handle(), /*y=*/output_data.opaque(), |
| 2125 | /*dyDesc=*/output_desc.data_handle(), |
| 2126 | /*dy=*/output_backprop_data.opaque(), nullptr, nullptr, |
| 2127 | /*dhyDesc=*/output_h_desc.handle(), |
nothing calls this directly
no test coverage detected