Python chainer.cuda.cupy() Examples

def mean_squared_error(x0, x1, ignore_nan=False):
    """Mean squared error function.

    This function computes mean squared error between two variables. The mean
    is taken over the minibatch. Note that the error is not scaled by 1/2.

    Args:
        x0 (:class:`~chainer.Variable` or :class:`numpy.ndarray` or \
        :class:`cupy.ndarray`): Input variable.
        x1 (:class:`~chainer.Variable` or :class:`numpy.ndarray` or \
        :class:`cupy.ndarray`): Input variable.
        ignore_nan (bool): If `True`, this function compute mean squared error
            ignoring NaNs. The arithmetic mean is the sum of the non-NaN
            elements along the axis divided by the number of whole elements.

    Returns:
        ~chainer.Variable:
            A variable holding an array representing the mean squared
            error of two inputs.
    """
    return MeanSquaredError(ignore_nan).apply((x0, x1))[0] 

def update(self, s, i):
        """Update decoder state

        Args:
            s (any): Current (hidden, cell) states.  If ``None`` is specified 
                     zero-vector is used.
            i (int): input label.
        Return:
            (~chainer.Variable) updated decoder state
        """
        if cuda.get_device_from_array(s[0].data).id >= 0:
            xp = cuda.cupy
        else:
            xp = np

        v = chainer.Variable(xp.array([i],dtype=np.int32))
        x = self.embed(v)
        if s is not None:
            hy, cy, dy = self.lstm(s[0], s[1], [x])
        else:
            hy, cy, dy = self.lstm(None, None, [x])

        return hy, cy, dy 

def update(self, s, i):
        """Update decoder state

        Args:
            s (any): Current (hidden, cell) states.  If ``None`` is specified 
                     zero-vector is used.
            i (int): input label.
        Return:
            (~chainer.Variable) updated decoder state
        """
        if cuda.get_device_from_array(s[0].data).id >= 0:
            xp = cuda.cupy
        else:
            xp = np

        v = chainer.Variable(xp.array([i],dtype=np.int32))
        x = self.embed(v)
        if s is not None:
            hy, cy, dy = self.lstm(s[0], s[1], [x])
        else:
            hy, cy, dy = self.lstm(None, None, [x])

        return hy, cy, dy 

def update(self, s, i):
        """Update decoder state

        Args:
            s (any): Current (hidden, cell) states.  If ``None`` is specified 
                     zero-vector is used.
            i (int): input label.
        Return:
            (~chainer.Variable) updated decoder state
        """
        if cuda.get_device_from_array(s[0].data).id >= 0:
            xp = cuda.cupy
        else:
            xp = np

        v = chainer.Variable(xp.array([i],dtype=np.int32))
        x = self.embed(v)
        if s is not None:
            hy, cy, dy = self.lstm(s[0], s[1], [x])
        else:
            hy, cy, dy = self.lstm(None, None, [x])

        return hy, cy, dy 

def update(self, s, i):
        """Update decoder state

        Args:
            s (any): Current (hidden, cell) states.  If ``None`` is specified 
                     zero-vector is used.
            i (int): input label.
        Return:
            (~chainer.Variable) updated decoder state
        """
        if cuda.get_device_from_array(s[0].data).id >= 0:
            xp = cuda.cupy
        else:
            xp = np

        v = chainer.Variable(xp.array([i],dtype=np.int32))
        x = self.embed(v)
        if s is not None:
            hy, cy, dy = self.lstm(s[0], s[1], [x])
        else:
            hy, cy, dy = self.lstm(None, None, [x])

        return hy, cy, dy 

def __init__(self, *, iterator, noise_iterator, optimizer_generator,
                 optimizer_critic, device=-1):

        if optimizer_generator.target.name is None:
            optimizer_generator.target.name = 'generator'

        if optimizer_critic.target.name is None:
            optimizer_critic.target.name = 'critic'

        iterators = {'main': iterator, 'z': noise_iterator}
        optimizers = {'generator': optimizer_generator,
                      'critic': optimizer_critic}

        super().__init__(iterators, optimizers, device=device)

        if device >= 0:
            cuda.get_device(device).use()
            [optimizer.target.to_gpu() for optimizer in optimizers.values()]

        self.xp = cuda.cupy if device >= 0 else np 

def __init__(self, in_channels, out_channels, ksize, stride=1, real=0, wscale=1.0):
        super(ConvolutionRBM, self).__init__(
            conv=L.Convolution2D(in_channels, out_channels, ksize, stride=stride, wscale=wscale),
        )

#        if gpu >= 0:
#            cuda.check_cuda_available()
#            xp = cuda.cupy # if gpu >= 0 else np
        self.conv.add_param("a", in_channels)  # dtype=xp.float32
        self.conv.a.data.fill(0.)
        self.in_channels = in_channels
        self.out_channels = out_channels
        self.ksize = ksize
        self.real = real

        self.rbm_train = False  # default value is false 

def __init__(self, args):
        xp = cuda.cupy if args.gpu >= 0 else np
        xp.random.seed(args.seed)
        Wh_data = xp.array([[[[1], [-1]]]], dtype='f')
        Ww_data = xp.array([[[[1, -1]]]], dtype='f')
        self.Wh = chainer.Variable(Wh_data)
        self.Ww = chainer.Variable(Ww_data)
        self.args = args

        self.load_model()
        self.create_dir()
        self.get_img_var()
        self.create_target()
        self.create_image_plane()
        self.prepare_optimizer()
        self.create_lr_schedule() 

def sample_h_given_v(self, v0_sample):
        """ get a sample of the hiddens by gibbs sampling
        :param v0_sample: Variable, see vis above
        :return:
        h1_mean:   Variable Matrix(batch_size, out_channels, image_height_out, image_width_out)
        h1_sample: Variable Matrix(batch_size, out_channels, image_height_out, image_width_out)
                   - actual sample for hidden units, populated by 0 or 1.
        """
        h1_mean = self.propup(v0_sample)
        xp = cuda.get_array_module(h1_mean.data)
        if xp == cuda.cupy:
            h1_sample = cuda.cupy.random.random_sample(size=h1_mean.data.shape)
            h1_sample[:] = h1_sample[:] < h1_mean.data[:]
        else:  # xp == np
            h1_sample = np.random.binomial(size=h1_mean.data.shape, n=1, p=h1_mean.data)
        return h1_mean, Variable(h1_sample.astype(xp.float32)) 

def backprop_check():
	xp = cuda.cupy if config.use_gpu else np
	duel = DDQN()

	state = xp.random.uniform(-1.0, 1.0, (2, config.rl_agent_history_length * config.ale_screen_channels, config.ale_scaled_screen_size[1], config.ale_scaled_screen_size[0])).astype(xp.float32)
	reward = [1, 0]
	action = [3, 4]
	episode_ends = [0, 0]
	next_state = xp.random.uniform(-1.0, 1.0, (2, config.rl_agent_history_length * config.ale_screen_channels, config.ale_scaled_screen_size[1], config.ale_scaled_screen_size[0])).astype(xp.float32)

	optimizer_conv = optimizers.Adam(alpha=config.rl_learning_rate, beta1=config.rl_gradient_momentum)
	optimizer_conv.setup(duel.conv)
	optimizer_fc = optimizers.Adam(alpha=config.rl_learning_rate, beta1=config.rl_gradient_momentum)
	optimizer_fc.setup(duel.fc)

	for i in xrange(10000):
		optimizer_conv.zero_grads()
		optimizer_fc.zero_grads()
		loss, _ = duel.forward_one_step(state, action, reward, next_state, episode_ends)
		loss.backward()
		optimizer_conv.update()
		optimizer_fc.update()
		print loss.data,
		print duel.conv.layer_2.W.data[0, 0, 0, 0],
		print duel.fc.layer_2.W.data[0, 0], 

def backward_gpu(self, inputs, grad_outputs):
        cupy = cuda.cupy
        x, t = inputs[:2]

        y = self.y
        gloss = grad_outputs[0]

        g_log_p = y
        g_log_p[cupy.arange(len(t)), cupy.maximum(t, 0)] -= 1

        g_log_p *= (t != self.ignore_label).reshape((len(t), 1))

        if self.reduce == 'mean':
            g_log_p *= gloss * self._coeff
        else:
            g_log_p *= gloss[:, None]

        ret = super(AdaptiveSoftmaxCrossEntropy, self).backward(
            inputs, (g_log_p, ))
        return ret 

def __call__(self, trainer):
        iteration = trainer.updater.iteration

        with cuda.get_device_from_id(trainer.updater.get_optimizer('main').target._device_id), chainer.using_config('train', False):
            self.xp = np if trainer.updater.get_optimizer('main').target._device_id < 0 else cuda.cupy
            image = self.xp.asarray(self.image)
            predictor = trainer.updater.get_optimizer('main').target.predictor
            predictions, rois, bboxes = predictor(image[self.xp.newaxis, ...])

            backprop_visualizations = []
            for visanchor in self.visualization_anchors:
                vis_target = predictor
                for target in visanchor:
                    vis_target = getattr(vis_target, target)
                backprop_visualizations.append(self.visual_backprop.perform_visual_backprop(vis_target))

            self.render_rois(predictions, rois, bboxes, iteration, self.image.copy(), backprop_vis=backprop_visualizations) 

def __init__(self, args):
        xp = cuda.cupy if args.gpu >= 0 else np
        xp.random.seed(args.seed)
        Wh_data = xp.array([[[[1], [-1]]]], dtype='f')
        Ww_data = xp.array([[[[1, -1]]]], dtype='f')
        self.Wh = chainer.Variable(Wh_data)
        self.Ww = chainer.Variable(Ww_data)
        self.args = args

        self.load_model()
        self.create_dir()
        self.get_img_var()
        self.create_target()
        self.create_image_plane()
        self.prepare_optimizer()
        self.create_lr_schedule() 

def backward_gpu(self, x, gy):
        xp = cuda.cupy
        gcol = conv.im2col_gpu(
            gy[0], self.kh, self.kw, self.sy, self.sx, self.ph, self.pw,
            cover_all=self.cover_all)

        gcol = gcol.transpose(0, 1, 4, 5, 2, 3)
        n, c, oy, ox, ky, kx = gcol.shape
        gcol = gcol.reshape((n, c, oy, ox, ky * kx))
        indexes = xp.asarray(self.indexes, dtype=numpy.int32)
        gx = xp.empty((n, c, oy, ox), dtype=x[0].dtype)
        xp.ElementwiseKernel(
            'int32 indexes, raw float32 gcol, int32 n, int32 c, int32 oy,'
            'int32 ox, int32 ky, int32 kx',
            'raw float32 gx',
            '''
            int ind_n = i / c / oy / ox;
            int ind_c = (i / oy / ox) % c;
            int ind_oy = (i / ox) % oy;
            int ind_ox = i % ox;
            int gcol_ky = indexes / kx;
            int gcol_kx = indexes % kx;
            float top_gx = gcol[ind_n * c * oy * ox * ky * kx + \
                                ind_c * oy * ox * ky * kx + \
                                ind_oy * ox * ky * kx + \
                                ind_ox * ky * kx + \
                                gcol_ky * kx + \
                                gcol_kx];
            gx[ind_n * c * oy * ox + \
               ind_c * oy * ox + \
               ind_oy * ox + \
               ind_ox] = top_gx;
            ''',
            'upsampling_2d_bwd')(indexes, gcol, n, c, oy, ox, ky, kx, gx)

        return gx, 

def forward(self, x_data, x_ref_data, y_data, train=True,
                 n_patches_per_image=32):

        xp = cuda.cupy

        if not isinstance(x_data, Variable):
            x = Variable(x_data)
        else:
            x = x_data
            x_data = x.data

        self.n_images = y_data.shape[0]
        self.n_patches = x_data.shape[0]
        self.n_patches_per_image = n_patches_per_image
        x_ref = Variable(x_ref_data)
       
        h = self.extract_features(x)
        self.h = h

        h_ref = self.extract_features(x_ref)

        h = F.concat((h-h_ref, h, h_ref))

        h_ = h # save intermediate features
        h = F.dropout(F.relu(self.fc1(h)), ratio=0.5)
        h = self.fc2(h)

        if self.top == "weighted":
            a = F.dropout(F.relu(self.fc1_a(h_)), ratio=0.5)
            a = F.relu(self.fc2_a(a))+0.000001
            t = Variable(y_data)
            self.weighted_loss(h, a, t)
        elif self.top == "patchwise":
            a = Variable(xp.ones_like(h.data))
            t = Variable(xp.repeat(y_data, n_patches_per_image))
            self.patchwise_loss(h, a, t)

        if train:
            return self.loss
        else:
            return self.loss, self.y 

def forward_gpu(self, x):
        xp = cuda.cupy
        n, c, h, w = x[0].shape
        if self.outh is None:
            self.outh = conv.get_deconv_outsize(
                h, self.kh, self.sy, self.ph, cover_all=self.cover_all)
        if self.outw is None:
            self.outw = conv.get_deconv_outsize(
                w, self.kw, self.sx, self.pw, cover_all=self.cover_all)
        up_y = xp.zeros((n, c, self.outh, self.outw), dtype=numpy.float32)
        up_y = conv.im2col_gpu(
            up_y, self.kh, self.kw, self.sy, self.sx, self.ph, self.pw,
            cover_all=self.cover_all)
        up_y = up_y.transpose(0, 1, 4, 5, 2, 3)
        n, c, oy, ox, ky, kx = up_y.shape
        indexes = xp.asarray(self.indexes, dtype=numpy.int32)
        xp.ElementwiseKernel(
            'int32 index, float32 x, int32 n, int32 c, int32 oy, int32 ox,'
            'int32 ky, int32 kx', 'raw float32 up_y',
            '''
            int yn = i / c / oy / ox;
            int yc = (i / oy / ox) % c;
            int yoy = (i / ox) % oy;
            int yox = i % ox;
            up_y[yn * c * oy * ox * ky * kx + \
              yc * oy * ox * ky * kx + \
              yoy * ox * ky * kx + \
              yox * ky * kx + \
              index] = x;
            ''',
            'upsampling_2d_fwd')(indexes, x[0], n, c, oy, ox, ky, kx, up_y)
        up_y = up_y.transpose(0, 1, 4, 5, 2, 3)
        up_y = conv.col2im_gpu(up_y, self.sy, self.sx, self.ph, self.pw,
                               self.outh, self.outw)
        return up_y, 

def forward_gpu(self, inputs):
        cupy = cuda.cupy
        x, t = inputs
        if chainer.is_debug():
            self._check_input_values(x, t)

        log_y = log_softmax._log_softmax(x, self.use_cudnn)
        if self.cache_score:
            self.y = cupy.exp(log_y)
        if self.class_weight is not None:
            shape = [1 if d != 1 else -1 for d in six.moves.range(x.ndim)]
            log_y *= cupy.broadcast_to(
                self.class_weight.reshape(shape), x.shape)
        if self.normalize:
            coeff = cupy.maximum(1, (t != self.ignore_label).sum())
        else:
            coeff = max(1, len(t))
        self._coeff = cupy.divide(1.0, coeff, dtype=x.dtype)

        log_y = cupy.rollaxis(log_y, 1, log_y.ndim)
        ret = cuda.reduce(
            'S t, raw T log_y, int32 n_channel, raw T coeff', 'T out',
            't == -1 ? T(0) : log_y[_j * n_channel + t]',
            'a + b', 'out = a * -coeff[0]', '0', 'crossent_fwd'
        )(t, log_y.reduced_view(), log_y.shape[-1], self._coeff)
        return ret, 

def backward_gpu(self, inputs, grad_outputs):
        cupy = cuda.cupy
        x, t = inputs
        if hasattr(self, 'y'):
            y = self.y
        else:
            y = log_softmax._log_softmax(x, self.use_cudnn)
            cupy.exp(y, out=y)
        gloss = grad_outputs[0]
        n_unit = t.size // len(t)
        coeff = gloss * self._coeff
        if self.class_weight is None:
            gx = cuda.elementwise(
                'T y, S t, raw T coeff, S n_channel, S n_unit',
                'T gx',
                '''
                    const int c = (i / n_unit % n_channel);
                    gx = (t == -1) ? 0 : (coeff[0] * (y - (c == t)));
                ''',
                'softmax_crossent_bwd')(
                    y, cupy.expand_dims(t, 1), coeff, x.shape[1], n_unit)
        else:
            gx = cuda.elementwise(
                'T y, raw T w, S t, raw T coeff, S n_channel, S n_unit',
                'T gx',
                '''
                    const int c = (i / n_unit % n_channel);
                    gx = t == -1 ? 0 : coeff[0] * (y - (c == t)) * w[t];
                ''',
                'softmax_crossent_bwd')(
                    y, self.class_weight, cupy.expand_dims(t, 1), coeff,
                    x.shape[1], n_unit)
        return gx, None 

def softmax_cross_entropy(
        x, t, use_cudnn=True, normalize=True, cache_score=True,
        class_weight=None):
    """Computes cross entropy loss for pre-softmax activations.

    Args:
        x (~chainer.Variable): Variable holding a multidimensional array whose
            element indicates unnormalized log probability: the first axis of
            the variable represents the number of samples, and the second axis
            represents the number of classes. While this function computes
            a usual softmax cross entropy if the number of dimensions is equal
            to 2, it computes a cross entropy of the replicated softmax if the
            number of dimensions is greater than 2.
        t (~chainer.Variable): Variable holding an int32 vector of ground truth
            labels. If ``t[i] == -1``, corresponding ``x[i]`` is ignored.
        normalize (bool): If ``True``, this function normalizes the cross
            entropy loss across all instances. If ``False``, it only
            normalizes along a batch size.
        cache_score (bool): When it is ``True``, the function stores result
            of forward computation to use it on backward computation. It
            reduces computational cost though consumes more memory.
        class_weight (~numpy.ndarray or ~cupy.ndarray): An array that contains
            constant weights that will be multiplied with the loss values along
            with the second dimension. The shape of this array should be
            ``(x.shape[1],)``.

    Returns:
        Variable: A variable holding a scalar array of the cross entropy loss.

    .. note::

       This function is differentiable only by ``x``.

    """
    return SoftmaxCrossEntropy(
        use_cudnn, normalize, cache_score, class_weight)(x, t) 

def check_call(
            self, gpu, label_key, args, kwargs, model_args, model_kwargs,
            metrics_fun, compute_metrics):
        init_kwargs = {'label_key': label_key}
        if metrics_fun is not None:
            init_kwargs['metrics_fun'] = metrics_fun
        link = Classifier(chainer.Link(), **init_kwargs)

        if gpu:
            xp = cuda.cupy
            link.to_gpu()
        else:
            xp = numpy

        link.compute_metrics = compute_metrics

        y = chainer.Variable(self.y)
        link.predictor = mock.MagicMock(return_value=y)

        loss = link(*args, **kwargs)
        link.predictor.assert_called_with(*model_args, **model_kwargs)

        assert hasattr(link, 'y')
        assert link.y is not None

        assert hasattr(link, 'loss')
        xp.testing.assert_allclose(link.loss.data, loss.data)

        assert hasattr(link, 'metrics')
        if compute_metrics:
            assert link.metrics is not None
        else:
            assert link.metrics is None 

def to_dense(self):
        """Returns a dense matrix format of this sparse matrix."""
        data = self.data
        if data.ndim == 1:
            shape = self.shape
        elif data.ndim == 2:
            shape = (data.shape[0], *self.shape)
        else:
            assert False

        xp = data.xp
        x = xp.zeros(shape, dtype=data.dtype)

        if data.size > 0:
            row = self.row
            col = self.col
            if xp is numpy:
                add_at = numpy.add.at
            elif xp is cuda.cupy:
                add_at = cuda.cupyx.scatter_add

            data = data.array
            if data.ndim == 1:
                _add_at(add_at, x, row, col, data)
            elif data.ndim == 2:
                for i in range(data.shape[0]):
                    _add_at(add_at, x[i], row[i], col[i], data[i])
            else:
                assert False
        return x 

def eps_greedy(self, state, exploration_rate):
		prop = np.random.uniform()
		q_max = None
		q_min = None
		if prop < exploration_rate:
			# Select a random action
			action_index = np.random.randint(0, len(config.ale_actions))
		else:
			# Select a greedy action
			state = Variable(state)
			if config.use_gpu:
				state.to_gpu()
			q = self.compute_q_variable(state, test=True)
			if config.use_gpu:
				action_index = cuda.to_cpu(cuda.cupy.argmax(q.data))
				q_max = cuda.to_cpu(cuda.cupy.max(q.data))
				q_min = cuda.to_cpu(cuda.cupy.min(q.data))
			else:
				action_index = np.argmax(q.data)
				q_max = np.max(q.data)
				q_min = np.min(q.data)

		action = self.get_action_with_index(action_index)
		# No-op
		self.no_op_count = self.no_op_count + 1 if action == 0 else 0
		if self.no_op_count > config.rl_no_op_max:
			no_op_index = np.argmin(np.asarray(config.ale_actions))
			actions_without_no_op = []
			for i in range(len(config.ale_actions)):
				if i == no_op_index:
					continue
				actions_without_no_op.append(config.ale_actions[i])
			action_index = np.random.randint(0, len(actions_without_no_op))
			action = actions_without_no_op[action_index]
			print "Reached no_op_max.", "New action:", action

		return action, q_max, q_min 

def output(self, h, t=None):
        Ws = [self.head] + [getattr(self, 'tail{}'.format(i))
                            for i in range(1, self.n_tails + 1)]
        Rs = [getattr(self, 'reduce{}'.format(i))
              for i in range(1, self.n_tails + 1)]
        cutoff = self.cutoff.data.astype('i').tolist()
        # An error happens to cupy when 0-dim array idx is directly used.

        output_all = t is None
        if output_all:
            t = self.xp.zeros((h.shape[0], ), 'i')
        return adaptive_softmax_output(
            h, t, Ws, Rs, cutoff, output_all=output_all) 

def forward_gpu(self, inputs):
        cupy = cuda.cupy
        x, t = inputs[:2]
        log_y = super(AdaptiveSoftmaxCrossEntropy, self).forward(inputs)[0]
        self.y = cupy.exp(log_y)

        if self.normalize:
            coeff = cupy.maximum(1, (t != self.ignore_label).sum())
        else:
            coeff = max(1, len(t))
        self._coeff = cupy.divide(1.0, coeff, dtype=x.dtype)

        log_y = cupy.rollaxis(log_y, 1, log_y.ndim)
        if self.reduce == 'mean':
            ret = cuda.reduce(
                'S t, raw T log_y, int32 n_channel, raw T coeff, '
                'S ignore_label',
                'T out',
                't == ignore_label ? T(0) : log_y[_j * n_channel + t]',
                'a + b', 'out = a * -coeff[0]', '0', 'crossent_fwd'
            )(t, log_y.reduced_view(), log_y.shape[-1],
              self._coeff, self.ignore_label)
        else:
            ret = cuda.elementwise(
                'S t, raw T log_y, int32 n_channel, T ignore', 'T out',
                '''
                if (t == ignore) {
                  out = 0;
                } else {
                  out = -log_y[i * n_channel + t];
                }
                ''',
                'softmax_crossent_no_reduce_fwd'
            )(t, log_y.reduced_view(), log_y.shape[-1], self.ignore_label)
            ret = ret.reshape(t.shape)
        return ret, 

def _check_class_weight_option(class_weight):
    if class_weight is not None:
        if class_weight.ndim != 1:
            raise ValueError('class_weight.ndim should be 1')
        if class_weight.dtype.kind != 'f':
            raise ValueError('The dtype of class_weight should be \'f\'')
        if isinstance(class_weight, variable.Variable):
            raise ValueError('class_weight should be a numpy.ndarray or '
                             'cupy.ndarray, not a chainer.Variable') 

def fit(self, x, y, valid_x, valid_y, img_siz, img_dim, test_x=None, test_y=None, callback=None):
        if self.device_id >= 0:
            with cuda.cupy.cuda.Device(self.device_id):
                return self.__fit(x, y, valid_x, valid_y, img_siz, img_dim, test_x, test_y, callback)
        else:
            return self.__fit(x, y, valid_x, valid_y, img_siz, img_dim, test_x, test_y, callback) 

def __init__(self, net, optimizer, epoch_num=100, batch_size=100, device_id=-1):
        self.net = net
        self.optimizer = optimizer
        self.epoch_num = epoch_num
        self.batch_size = batch_size
        self.device_id = device_id
        if device_id >= 0:
            self.xp = cuda.cupy
            self.net.to_gpu(device_id)
        else:
            self.xp = np 

def __call__(self, rule, param):
        p, g = param.data, param.grad
        if p is None or g is None:
            return
        with chainer.using_device(param.device):
            xp = param.device.xp
            sign = xp.sign(p)
            if xp is cuda.cupy:
                kernel = cuda.elementwise(
                    'T s, T decay', 'T g', 'g += decay * s', 'lasso')
                kernel(sign, self.rate, g)
            else:
                g += self.rate * sign 

def __call__(self, rule, param):
        p, g = param.data, param.grad
        if p is None or g is None:
            return
        with chainer.using_device(param.device):
            rate = self.rate
            if param._loss_scale is not None:
                rate *= param._loss_scale
            if param.device.xp is cuda.cupy:
                kernel = cuda.elementwise(
                    'T p, T decay', 'T g', 'g += decay * p', 'weight_decay')
                kernel(p, rate, g)
            else:
                g += rate * p 

def test_unstride_backward_gpu(self):
        self.check_unstride_backward(cuda.cupy)