| 1263 | } |
| 1264 | |
| 1265 | void tensor_conv::get_gradient_for_data ( |
| 1266 | const bool add_to_output, |
| 1267 | const tensor& gradient_input, |
| 1268 | const tensor& filters, |
| 1269 | tensor& data_gradient |
| 1270 | ) |
| 1271 | { |
| 1272 | const float alpha = 1; |
| 1273 | const float beta = add_to_output ? 1 : 0; |
| 1274 | |
| 1275 | // Since cudnnConvolutionBackwardData() is an asynchronous call, we need to hold a |
| 1276 | // reference to the workspace buffer so we can be sure it isn't reallocated |
| 1277 | // while the function is still executing on the device. But each time we come |
| 1278 | // here, we make sure to grab the latest workspace buffer so that, globally, we |
| 1279 | // minimize the number of such buffers. |
| 1280 | backward_data_workspace = device_global_buffer(backward_data_workspace_size_in_bytes); |
| 1281 | |
| 1282 | |
| 1283 | CHECK_CUDNN(cudnnConvolutionBackwardData(context(), |
| 1284 | &alpha, |
| 1285 | (const cudnnFilterDescriptor_t)filter_handle, |
| 1286 | filters.device(), |
| 1287 | descriptor(gradient_input), |
| 1288 | gradient_input.device(), |
| 1289 | (const cudnnConvolutionDescriptor_t)conv_handle, |
| 1290 | (cudnnConvolutionBwdDataAlgo_t)backward_data_algo, |
| 1291 | backward_data_workspace, |
| 1292 | backward_data_workspace_size_in_bytes, |
| 1293 | &beta, |
| 1294 | descriptor(data_gradient), |
| 1295 | data_gradient.device())); |
| 1296 | } |
| 1297 | |
| 1298 | void tensor_conv:: |
| 1299 | get_gradient_for_filters ( |