def forward_gpu(self, x): n, out_c, out_h, out_w = x[0].shape c = self.W.shape[1] h = get_deconv_outsize(out_h, self.kh, self.sy, self.ph) w = get_deconv_outsize(out_w, self.kw, self.sx, self.pw) if cudnn.enabled and self.use_cudnn: handle = cudnn.get_default_handle() x_desc = cudnn.get_tensor_desc(x[0], out_h, out_w) y = cuda.empty((n, c, h, w), dtype=numpy.float32) y_desc = cudnn.get_tensor_desc(y, h, w) self.filter_desc = cudnn.get_filter4d_desc(self.W) self.conv_desc = cudnn.get_conv2d_desc( (self.ph, self.pw), (self.sy, self.sx)) if self.b is not None: self.bias_desc = cudnn.get_conv_bias_desc(self.b) libcudnn.cudnnConvolutionBackwardData( handle, 1, self.filter_desc.value, cudnn.get_ptr(self.W), x_desc.value, cudnn.get_ptr(x[0]), self.conv_desc.value, 0, y_desc.value, cudnn.get_ptr(y)) if self.b is not None: libcudnn.cudnnAddTensor( handle, libcudnn.cudnnAddMode['CUDNN_ADD_SAME_C'], 1, self.bias_desc.value, cudnn.get_ptr(self.b), 1, y_desc.value, cudnn.get_ptr(y)) else: handle = cuda.get_cublas_handle() # TODO(beam2d): Use streams W_mat = self.W.reshape(out_c, c * self.kh * self.kw) x_mats = x[0].reshape(n, out_c, out_h * out_w) gcol = cuda.empty((n, c, self.kh, self.kw, out_h, out_w), dtype=numpy.float32) gcol_mats = gcol.reshape(n, c * self.kh * self.kw, out_h * out_w) for i in moves.range(n): cuda.culinalg.dot(W_mat, x_mats[i], transa='T', handle=handle, out=gcol_mats[i]) y = conv.col2im_gpu( gcol, self.sy, self.sx, self.ph, self.pw, h, w) # TODO(beam2d): Support unshared bias if self.b is not None: cuda.elementwise( 'float* y, const float* b, int c, int hw', 'y[i] += b[i / hw % c]', 'conv_bias_fwd')(y, self.b, c, h * w) return y,
def forward_gpu(self, x): n, c, h, w = x[0].shape out_h = conv.get_conv_outsize(h, self.kh, self.sy, self.ph) out_w = conv.get_conv_outsize(w, self.kw, self.sx, self.pw) out_c = self.W.shape[0] y = cuda.empty((n, out_c, out_h, out_w), dtype=numpy.float32) if cudnn.enabled and self.use_cudnn: handle = cudnn.get_default_handle() x_desc = cudnn.get_tensor_desc(x[0], h, w) y_desc = cudnn.get_tensor_desc(y, out_h, out_w) self.filter_desc = cudnn.get_filter4d_desc(self.W) self.conv_desc = cudnn.get_conv2d_desc( (self.ph, self.pw), (self.sy, self.sx)) if self.b is not None: self.bias_desc = cudnn.get_conv_bias_desc(self.b) algo = libcudnn.cudnnGetConvolutionForwardAlgorithm( handle, x_desc.value, self.filter_desc.value, self.conv_desc.value, y_desc.value, _fwd_pref, self.max_workspace_size) workspace_size = libcudnn.cudnnGetConvolutionForwardWorkspaceSize( handle, x_desc.value, self.filter_desc.value, self.conv_desc.value, y_desc.value, algo).value workspace = cuda.empty( (max(workspace_size // 4, 1),), dtype=numpy.float32) libcudnn.cudnnConvolutionForward( handle, 1, x_desc.value, cudnn.get_ptr(x[0]), self.filter_desc.value, cudnn.get_ptr(self.W), self.conv_desc.value, algo, cudnn.get_ptr( workspace), workspace_size, 0, y_desc.value, cudnn.get_ptr(y)) # TODO(beam2d): Support unshared bias if self.b is not None: libcudnn.cudnnAddTensor( handle, libcudnn.cudnnAddMode['CUDNN_ADD_SAME_C'], 1, self.bias_desc.value, cudnn.get_ptr(self.b), 1, y_desc.value, cudnn.get_ptr(y)) else: # Implementation using im2col self.col = conv.im2col_gpu( x[0], self.kh, self.kw, self.sy, self.sx, self.ph, self.pw) # TODO(beam2d): Use streams handle = cuda.get_cublas_handle() W_mat = self.W.reshape(out_c, c * self.kh * self.kw) col_mats = self.col.reshape( n, c * self.kh * self.kw, out_h * out_w) y_mats = y.reshape(n, out_c, out_h * out_w) for i in moves.range(n): cuda.culinalg.dot(W_mat, col_mats[i], handle=handle, out=y_mats[i]) # TODO(beam2d): Support unshared bias if self.b is not None: cuda.elementwise( 'float* y, const float* b, int c, int hw', 'y[i] += b[i / hw % c]', 'conv_bias_fwd')(y, self.b, out_c, out_h * out_w) return y,