void ConvBC01CuDNN<T>::bprop(const T* imgs, const T* filters,
const T *convout_d, T *imgs_d, T *filters_d) {
if (filters_d) {
CUDNN_CHECK(cudnnConvolutionBackwardFilter(
CUDNN::handle(), &CUDNN::one, imgs_desc, imgs, convout_desc, convout_d,
conv_desc, &CUDNN::zero, filters_desc, filters_d
));
}
if (imgs_d) {
CUDNN_CHECK(cudnnConvolutionBackwardData(
CUDNN::handle(), &CUDNN::one, filters_desc, filters, convout_desc,
convout_d, conv_desc, &CUDNN::zero, imgs_desc, imgs_d
));
}
}
int
APPLY_SPECIFIC(conv_gw)(CudaNdarray *input, CudaNdarray *output,
cudnnConvolutionDescriptor_t desc,
int h, int w,
CudaNdarray **kerns) {
cudnnStatus_t err = CUDNN_STATUS_SUCCESS;
if (c_set_tensor4d(input, APPLY_SPECIFIC(input)) == -1)
return 1;
if (c_set_tensor4d(output, APPLY_SPECIFIC(output)) == -1)
return 1;
{
int out_dims[4];
out_dims[0] = CudaNdarray_HOST_DIMS(output)[1];
out_dims[1] = CudaNdarray_HOST_DIMS(input)[1];
out_dims[2] = h;
out_dims[3] = w;
if (CudaNdarray_prep_output(kerns, 4, out_dims) != 0) {
return 1;
}
}
if (c_set_filter(*kerns, APPLY_SPECIFIC(kerns)) == -1)
return 1;
{
const float alpha = 1;
const float beta = 0;
err = cudnnConvolutionBackwardFilter(
_handle,
(void *)&alpha,
APPLY_SPECIFIC(input), CudaNdarray_DEV_DATA(input),
APPLY_SPECIFIC(output), CudaNdarray_DEV_DATA(output),
desc,
(void *)&beta,
APPLY_SPECIFIC(kerns), CudaNdarray_DEV_DATA(*kerns));
}
if (err != CUDNN_STATUS_SUCCESS) {
PyErr_Format(PyExc_RuntimeError, "GpuDnnConvGradW: error doing operation: %s",
cudnnGetErrorString(err));
return 1;
}
return 0;
}
void ConvBC01CuDNN<T>::bprop(const T* imgs, const T* filters,
const T *convout_d, T *imgs_d, T *filters_d) {
void *workspace = NULL;
if (workspace_size > 0) {
workspace = CUDA::buffer(workspace_size);
}
if (filters_d) {
CUDNN_CHECK(cudnnConvolutionBackwardFilter(
CUDNN::handle(), &CUDNN::one, imgs_desc, imgs, convout_desc, convout_d,
conv_desc, bwd_filters_algo, workspace, workspace_size, &CUDNN::zero,
filters_desc, filters_d
));
}
if (imgs_d) {
CUDNN_CHECK(cudnnConvolutionBackwardData(
CUDNN::handle(), &CUDNN::one, filters_desc, filters, convout_desc,
convout_d, conv_desc, bwd_imgs_algo, workspace, workspace_size,
&CUDNN::zero, imgs_desc, imgs_d
));
}
}
int
APPLY_SPECIFIC(conv_gw)(CudaNdarray *input, CudaNdarray *output,
CudaNdarray *km, cudnnConvolutionDescriptor_t desc,
float alpha, float beta, CudaNdarray **kerns) {
cudnnStatus_t err = CUDNN_STATUS_SUCCESS;
if (c_set_tensor5d(input, APPLY_SPECIFIC(input)) == -1)
return 1;
if (c_set_tensor5d(output, APPLY_SPECIFIC(output)) == -1)
return 1;
#ifdef CONV_INPLACE
Py_XDECREF(*kerns);
*kerns = km;
Py_INCREF(*kerns);
#else
if (CudaNdarray_prep_output(kerns, 5, CudaNdarray_HOST_DIMS(km)) != 0)
return 1;
if (beta != 0.0 && CudaNdarray_CopyFromCudaNdarray(*kerns, km))
return 1;
#endif
if (c_set_filter5d(*kerns, APPLY_SPECIFIC(kerns)) == -1)
return 1;
err = cudnnConvolutionBackwardFilter(
_handle,
(void *)&alpha,
APPLY_SPECIFIC(input), CudaNdarray_DEV_DATA(input),
APPLY_SPECIFIC(output), CudaNdarray_DEV_DATA(output),
desc,
(void *)&beta,
APPLY_SPECIFIC(kerns), CudaNdarray_DEV_DATA(*kerns));
if (err != CUDNN_STATUS_SUCCESS) {
PyErr_Format(PyExc_RuntimeError, "GpuDnnConvGradW: error doing operation: %s",
cudnnGetErrorString(err));
return 1;
}
return 0;
}
void convolution_layer_updater_cuda::enqueue_update_weights(
unsigned int offset_input_entry_id,
cudaStream_t stream_id,
const std::vector<cuda_linear_buffer_device_smart_ptr>& gradient,
const std::vector<cuda_linear_buffer_device_smart_ptr>& data_custom,
const std::vector<const_cuda_linear_buffer_device_smart_ptr>& schema_data,
cuda_linear_buffer_device_smart_ptr output_errors_buffer,
const_cuda_linear_buffer_device_smart_ptr input_neurons_buffer,
const std::vector<cuda_linear_buffer_device_smart_ptr>& additional_buffers,
std::vector<cuda_memobject_smart_ptr>& dynamic_memobjects,
unsigned int entry_count,
bool force_deterministic)
{
cudnn_safe_call(cudnnSetStream(cuda_config->get_cudnn_handle(), stream_id));
cudnn_safe_call(cudnnSetTensor4dDescriptor(
input_data_desc,
CUDNN_TENSOR_NCHW,
CUDNN_DATA_FLOAT,
entry_count,
input_configuration_specific.feature_map_count,
(input_configuration_specific.dimension_sizes.size() > 1) ? input_configuration_specific.dimension_sizes[1] : 1,
input_configuration_specific.dimension_sizes[0]));
cudnn_safe_call(cudnnSetTensor4dDescriptor(
output_data_desc,
CUDNN_TENSOR_NCHW,
CUDNN_DATA_FLOAT,
entry_count,
output_configuration_specific.feature_map_count,
(output_configuration_specific.dimension_sizes.size() > 1) ? output_configuration_specific.dimension_sizes[1] : 1,
output_configuration_specific.dimension_sizes[0]));
{
float alpha = 1.0F;
float beta = 1.0F;
cudnn_safe_call(cudnnConvolutionBackwardFilter(
cuda_config->get_cudnn_handle(),
&alpha,
input_data_desc,
(const float *)(*input_neurons_buffer) + input_elem_count_per_entry * offset_input_entry_id,
output_data_desc,
*output_errors_buffer,
convolution_desc,
&beta,
weights_desc,
*gradient[0]));
}
{
float alpha = 1.0F;
float beta = 1.0F;
cudnn_safe_call(cudnnConvolutionBackwardBias(
cuda_config->get_cudnn_handle(),
&alpha,
output_data_desc,
*output_errors_buffer,
&beta,
bias_desc,
*gradient[1]));
}
}
int
APPLY_SPECIFIC(conv_gw)(CudaNdarray *input, CudaNdarray *output,
CudaNdarray *km, cudnnConvolutionDescriptor_t desc,
float alpha, float beta, CudaNdarray **kerns) {
cudnnStatus_t err = CUDNN_STATUS_SUCCESS;
if (CudaNdarray_HOST_DIMS(input)[1] != CudaNdarray_HOST_DIMS(km)[1]) {
PyErr_SetString(PyExc_ValueError,
"GpuDnnConv images and kernel must have the same stack size\n");
return 1;
}
if (c_set_tensorNd(input, APPLY_SPECIFIC(input)) == -1)
return 1;
if (c_set_tensorNd(output, APPLY_SPECIFIC(output)) == -1)
return 1;
int nb_dim = CudaNdarray_NDIM(output);
#ifdef CONV_INPLACE
Py_XDECREF(*kerns);
*kerns = km;
Py_INCREF(*kerns);
#else
if (CudaNdarray_prep_output(kerns, nb_dim, CudaNdarray_HOST_DIMS(km)) != 0)
return 1;
if (beta != 0.0 && CudaNdarray_CopyFromCudaNdarray(*kerns, km))
return 1;
#endif
if (c_set_filterNd(*kerns, APPLY_SPECIFIC(kerns)) == -1)
return 1;
#if defined(CUDNN_VERSION) && CUDNN_VERSION >= 3000
{
size_t worksize;
void *workspace;
cudnnConvolutionBwdFilterAlgo_t chosen_algo;
if (CHOOSE_ALGO)
{
// A new convolution implementation should be selected, based either on
// timing or heuristics, if in one of the two following cases :
// - The implementation should only be chosen during the first execution
// of an apply node and this is the first execution of the apply node.
// - The implementation should be chosen as often as necessary and the
// shapes of the inputs differ from the last time an implementation
// was chosen.
bool reuse_previous_algo;
if (CHOOSE_ALGO_ONCE)
{
// Only choose a new implementation of none has been chosen before.
reuse_previous_algo = APPLY_SPECIFIC(previous_algo_set);
}
else
{
// Reuse the previous implementation if the the kernels and the outputs
// have the same shapes as they had when the previous implementation
// was selected
bool same_shapes = true;
for (int i = 0; (i < nb_dim) && same_shapes; i++)
{
same_shapes &= (CudaNdarray_HOST_DIMS(input)[i] ==
APPLY_SPECIFIC(previous_input_shape)[i]);
same_shapes &= (CudaNdarray_HOST_DIMS(output)[i] ==
APPLY_SPECIFIC(previous_output_shape)[i]);
}
reuse_previous_algo = same_shapes;
}
// If the previously choosen implementation can't be reused, select a
// new one based on the shapes of the current inputs
if (!reuse_previous_algo)
{
// Obtain a convolution algorithm appropriate for the input and output
// shapes. Either by choosing one according to heuristics or by making
// CuDNN time every implementation and choose the best one.
if (CHOOSE_ALGO_TIME)
{
// Time the different implementations to choose the best one
int requestedCount = 1;
int count;
cudnnConvolutionBwdFilterAlgoPerf_t choosen_algo_perf;
err = cudnnFindConvolutionBackwardFilterAlgorithm(_handle,
APPLY_SPECIFIC(input),
APPLY_SPECIFIC(output),
desc,
APPLY_SPECIFIC(kerns),
requestedCount,
&count,
&choosen_algo_perf);
if (err != CUDNN_STATUS_SUCCESS) {
PyErr_Format(PyExc_RuntimeError,
"GpuDnnConvGradW: error selecting convolution algo: "
"%s", cudnnGetErrorString(err));
return 1;
}
chosen_algo = choosen_algo_perf.algo;
}
else
{
// Choose the convolution implementation using heuristics based on the
// shapes of the inputs and the amount of memory available.
// Get the amount of available memory
size_t free = 0, total = 0;
cudaError_t err2 = cudaMemGetInfo(&free, &total);
if (err2 != cudaSuccess){
cudaGetLastError();
fprintf(stderr,
"Error when trying to find the memory information"
" on the GPU: %s\n", cudaGetErrorString(err2));
return 1;
}
// Use heuristics to choose the implementation
err = cudnnGetConvolutionBackwardFilterAlgorithm(_handle,
APPLY_SPECIFIC(input),
APPLY_SPECIFIC(output),
desc,
APPLY_SPECIFIC(kerns),
CUDNN_CONVOLUTION_BWD_FILTER_SPECIFY_WORKSPACE_LIMIT,
free,
&chosen_algo);
if (err != CUDNN_STATUS_SUCCESS) {
PyErr_Format(PyExc_RuntimeError,
"GpuDnnConvGradW: error selecting convolution algo: %s",
cudnnGetErrorString(err));
return 1;
}
}
// Store the shapes of the inputs and kernels as well as the chosen
// algorithm for future use.
APPLY_SPECIFIC(previous_bwd_f_algo) = chosen_algo;
APPLY_SPECIFIC(previous_algo_set) = true;
for (int i = 0; i < nb_dim; i++)
{
APPLY_SPECIFIC(previous_input_shape)[i] =
CudaNdarray_HOST_DIMS(input)[i];
APPLY_SPECIFIC(previous_output_shape)[i] =
CudaNdarray_HOST_DIMS(output)[i];
}
}
else
{
// Reuse the previously chosen convlution implementation
chosen_algo = APPLY_SPECIFIC(previous_bwd_f_algo);
}
}
else
{
chosen_algo = CONV_ALGO;
}
// The FFT implementation (only in v3 and onward) does not support strides,
// 1x1 filters or inputs with a spatial dimension larger than 1024.
// If the chosen implementation is FFT, validate that it can be used
// on the current data and default on a safe implementation if it
// can't.
if (chosen_algo == CUDNN_CONVOLUTION_BWD_FILTER_ALGO_FFT && nb_dim == 4)
{
// Extract the properties of the convolution descriptor
int pad_h, pad_w, stride_v, stride_h, upscale_x, upscale_y;
cudnnConvolutionMode_t mode;
err = cudnnGetConvolution2dDescriptor(desc, &pad_h, &pad_w,
&stride_v, &stride_h,
&upscale_x, &upscale_y,
&mode);
if (err != CUDNN_STATUS_SUCCESS) {
PyErr_Format(PyExc_RuntimeError,
"GpuDnnConvGradW: error getting convolution properties: %s",
cudnnGetErrorString(err));
return 1;
}
// Extract the spatial size of the filters
int filter_h = CudaNdarray_HOST_DIMS(*kerns)[2];
int filter_w = CudaNdarray_HOST_DIMS(*kerns)[3];
// Extract the spatial size of the input
int input_h = CudaNdarray_HOST_DIMS(input)[2];
int input_w = CudaNdarray_HOST_DIMS(input)[3];
// Ensure that the selected implementation supports the requested
// convolution. Fall back to a safe implementation otherwise.
if (stride_v != 1 || stride_h != 1 || input_h > 1024 ||
input_w > 1024 || (filter_h == 1 && filter_w == 1))
{
chosen_algo = CUDNN_CONVOLUTION_BWD_FILTER_ALGO_0;
}
}
// Infer required workspace size from the chosen implementation
err = cudnnGetConvolutionBackwardFilterWorkspaceSize(_handle,
APPLY_SPECIFIC(input),
APPLY_SPECIFIC(output),
desc,
APPLY_SPECIFIC(kerns),
chosen_algo,
&worksize);
if (err != CUDNN_STATUS_SUCCESS) {
PyErr_Format(PyExc_RuntimeError,
"GpuDnnConvGradW: error getting worksize: %s",
cudnnGetErrorString(err));
return 1;
}
// Allocate workspace for the convolution
workspace = get_work_mem(worksize);
if (workspace == NULL && worksize != 0)
return 1;
// Perform the convolution
err = cudnnConvolutionBackwardFilter_v3(
_handle,
(void *)&alpha,
APPLY_SPECIFIC(input), CudaNdarray_DEV_DATA(input),
APPLY_SPECIFIC(output), CudaNdarray_DEV_DATA(output),
desc,
chosen_algo,
workspace, worksize,
(void *)&beta,
APPLY_SPECIFIC(kerns), CudaNdarray_DEV_DATA(*kerns));
}
#else
err = cudnnConvolutionBackwardFilter(
_handle,
(void *)&alpha,
APPLY_SPECIFIC(input), CudaNdarray_DEV_DATA(input),
APPLY_SPECIFIC(output), CudaNdarray_DEV_DATA(output),
desc,
(void *)&beta,
APPLY_SPECIFIC(kerns), CudaNdarray_DEV_DATA(*kerns));
#endif
if (err != CUDNN_STATUS_SUCCESS) {
PyErr_Format(PyExc_RuntimeError, "GpuDnnConvGradW: error doing operation: %s",
cudnnGetErrorString(err));
return 1;
}
return 0;
}
static PyObject *conv_dfilter_buffers(PyObject *self, PyObject *args) {
cudaError_t err;
cudnnStatus_t status;
int PAD, gpu_ind, filters_ind, imgs_ind, conv_out_ind, out_ind, stream_ind;
if (!PyArg_ParseTuple(args, "iiiiiii", &filters_ind, &imgs_ind, &conv_out_ind, &out_ind, &PAD, &stream_ind, &gpu_ind))
return NULL;
if(filters_ind >= N_BUFFERS || filters_ind < 0 || imgs_ind >= N_BUFFERS || imgs_ind < 0 ||
conv_out_ind >= N_BUFFERS || conv_out_ind < 0 || out_ind >= N_BUFFERS || out_ind < 0){
printf("invalid buffer index\n");
return NULL;
}
if(gpu_ind < 0 || gpu_ind > N_GPUS){
printf("invalid gpu index %i\n", gpu_ind);
return NULL;
}
if(stream_ind < 0 || stream_ind > N_ALT_STREAMS){
printf("invalid stream index %i\n", stream_ind);
return NULL;
}
if(data_buffers[gpu_ind][filters_ind] == NULL || data_buffers[gpu_ind][imgs_ind] == NULL ||
data_buffers[gpu_ind][conv_out_ind] == NULL){
printf("one or more buffers not initialized on this gpu\n");
return NULL;
}
if(filter_flags[gpu_ind][filters_ind] == 0 || filter_flags[gpu_ind][imgs_ind] == 1 ||
filter_flags[gpu_ind][conv_out_ind] == 1){
printf("one or more buffers was not initialized correctly, filters when should be tensor or vice versa\n");
return NULL;
}
cudaSetDevice(gpu_ind); CHECK_CUDA_ERR
cudaStreamSynchronize(streams[gpu_ind]); // make sure the inputs are in the buffers first
cudnnSetStream(handle, alt_streams[gpu_ind][stream_ind]);
int n_filters = data_dims[0][gpu_ind][filters_ind];
int n_channels = data_dims[1][gpu_ind][filters_ind];
int filter_sz = data_dims[2][gpu_ind][filters_ind];
if(data_buffers[gpu_ind][out_ind] == NULL){ // allocate output
status = cudnnCreateFilterDescriptor(&desc_filters[gpu_ind][out_ind]); ERR_CHECK
status = cudnnSetFilterDescriptor(desc_filters[gpu_ind][out_ind], dataType, n_filters, n_channels, filter_sz, filter_sz); ERR_CHECK
err = cudaMalloc((void**) &data_buffers[gpu_ind][out_ind], n_filters*n_channels*filter_sz*filter_sz * DATA_TYPE_SZ); MALLOC_ERR_CHECK
data_dims[0][gpu_ind][out_ind] = n_filters;
data_dims[1][gpu_ind][out_ind] = n_channels;
data_dims[2][gpu_ind][out_ind] = filter_sz;
data_dims[3][gpu_ind][out_ind] = filter_sz;
filter_flags[gpu_ind][out_ind] = 1;
}else if(filter_flags[gpu_ind][out_ind] == 0 || data_dims[0][gpu_ind][out_ind] != n_filters ||
data_dims[1][gpu_ind][out_ind] != n_channels || data_dims[2][gpu_ind][out_ind] != filter_sz ||
data_dims[3][gpu_ind][out_ind] != filter_sz){ // make sure output buffer is of correct size
printf("output buffer size is not matching output of this function and/or initialized as a tensor, %s %i\n", __FILE__, __LINE__);
return NULL;
}
//---------------------------------------
// Set decriptors
//---------------------------------------
status = cudnnSetConvolutionDescriptor(convDesc, desc_buffers[gpu_ind][imgs_ind], desc_filters[gpu_ind][out_ind], PAD, PAD, 1, 1, 1, 1, CUDNN_CROSS_CORRELATION); ERR_CHECK
//---------------------------------------
// Query output layout
//---------------------------------------
int n_imgs_out, n_filters_out, conv_out_sz_x, conv_out_sz_y;
status = cudnnGetOutputTensor4dDim(convDesc, CUDNN_CONVOLUTION_FWD, &n_imgs_out, &n_filters_out, &conv_out_sz_x, &conv_out_sz_y); ERR_CHECK
//--------------------------------------
// set filter and image values
//--------------------------------------
if(n_imgs_out*n_filters_out*conv_out_sz_x*conv_out_sz_x != data_dims[0][gpu_ind][conv_out_ind]*data_dims[1][gpu_ind][conv_out_ind]*
data_dims[2][gpu_ind][conv_out_ind]*data_dims[3][gpu_ind][conv_out_ind]){
printf("predicted conv output not matching given input %s %i\n", __FILE__, __LINE__);
printf("%i %i\n", n_imgs_out*n_filters_out*conv_out_sz_x*conv_out_sz_x, data_dims[0][gpu_ind][conv_out_ind]*data_dims[1][gpu_ind][conv_out_ind]*
data_dims[2][gpu_ind][conv_out_ind]*data_dims[3][gpu_ind][conv_out_ind]);
printf("%i %i\n", n_imgs_out, data_dims[0][gpu_ind][conv_out_ind]);
printf("%i %i\n", n_filters_out, data_dims[1][gpu_ind][conv_out_ind]);
printf("%i %i\n", conv_out_sz_x, data_dims[2][gpu_ind][conv_out_ind]);
printf("%i %i\n", conv_out_sz_y, data_dims[3][gpu_ind][conv_out_ind]);
//return NULL;
}
//--------------------------------------
// Convolution
//--------------------------------------
status = cudnnConvolutionBackwardFilter(handle, desc_buffers[gpu_ind][imgs_ind], data_buffers[gpu_ind][imgs_ind],
desc_buffers[gpu_ind][conv_out_ind], data_buffers[gpu_ind][conv_out_ind], convDesc,
desc_filters[gpu_ind][out_ind], data_buffers[gpu_ind][out_ind], CUDNN_RESULT_NO_ACCUMULATE); ERR_CHECK
cudnnSetStream(handle, NULL);
cudaSetDevice(0); CHECK_CUDA_ERR
Py_INCREF(Py_None);
return Py_None;
}
void convolution_layer_updater_cuda::enqueue_backward_weights_propagation(
cudaStream_t stream_id,
const std::vector<cuda_linear_buffer_device::const_ptr>& schema_data,
const std::vector<cuda_linear_buffer_device::ptr>& gradient,
const std::vector<cuda_linear_buffer_device::const_ptr>& data_custom,
const std::vector<cuda_linear_buffer_device::const_ptr>& input_neurons_buffers,
cuda_linear_buffer_device::const_ptr output_errors_buffer,
const std::vector<cuda_linear_buffer_device::const_ptr>& persistent_working_data,
cuda_linear_buffer_device::ptr temporary_working_fixed_buffer,
cuda_linear_buffer_device::ptr temporary_working_per_entry_buffer,
cuda_linear_buffer_device::const_ptr temporary_fixed_buffer,
cuda_linear_buffer_device::const_ptr temporary_per_entry_buffer,
unsigned int entry_count)
{
cudnn_safe_call(cudnnSetStream(cuda_config->get_cudnn_handle(), stream_id));
cudnn_util::set_tensor_descriptor(
input_data_desc,
input_configuration_specific_list[0],
entry_count);
cudnn_util::set_tensor_descriptor(
output_data_desc,
output_configuration_specific,
entry_count);
{
void * workspace = 0;
size_t workspace_size = 0;
if (temporary_working_fixed_buffer)
{
workspace = *temporary_working_fixed_buffer;
workspace_size = temporary_working_fixed_buffer->get_size();
}
cudnnConvolutionBwdFilterAlgo_t algo = cuda_config->cudnn_find_convolution_backward_weights_algo(
input_data_desc,
weights_desc,
convolution_desc,
output_data_desc,
*input_neurons_buffers[0],
*output_errors_buffer,
(unsigned char *)workspace,
(unsigned char *)workspace + update_weights_find_algo_working_buffer_size,
workspace_size - update_weights_find_algo_working_buffer_size);
float alpha = 1.0F;
float beta = 1.0F;
cudnn_safe_call(cudnnConvolutionBackwardFilter(
cuda_config->get_cudnn_handle(),
&alpha,
input_data_desc,
*input_neurons_buffers[0],
output_data_desc,
*output_errors_buffer,
convolution_desc,
algo,
workspace,
workspace_size,
&beta,
weights_desc,
*gradient[0]));
}
if (bias)
{
float alpha = 1.0F;
float beta = 1.0F;
cudnn_safe_call(cudnnConvolutionBackwardBias(
cuda_config->get_cudnn_handle(),
&alpha,
output_data_desc,
*output_errors_buffer,
&beta,
bias_desc,
*gradient[1]));
}
}