forked from pytorch/pytorch
-
Notifications
You must be signed in to change notification settings - Fork 0
/
L1Cost.cu
44 lines (35 loc) · 1.45 KB
/
L1Cost.cu
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
#ifndef THC_GENERIC_FILE
#define THC_GENERIC_FILE "THCUNN/generic/L1Cost.cu"
#else
void THNN_(L1Cost_updateOutput)(
THCState *state,
THCTensor *input,
THCTensor *output)
{
THCUNN_check_dim_size(state, output, 1, 0, 1);
THCUNN_assertSameGPU(state, 1, input);
accreal sum;
ptrdiff_t size = THCTensor_(nElement)(state, input);
input = THCTensor_(newContiguous)(state, input);
thrust::device_ptr<scalar_t> input_data(THCTensor_(data)(state, input));
sum = thrust::transform_reduce(input_data, input_data+size, l1cost_functor<scalar_t, accreal>(), accreal(0), thrust::plus<accreal>());
THCTensor_(free)(state, input);
THCTensor_(set1d)(state, output, 0, ScalarConvert<accreal, scalar_t>::to(sum));
}
void THNN_(L1Cost_updateGradInput)(
THCState *state,
THCTensor *input,
THCTensor *gradOutput,
THCTensor *gradInput)
{
THCUNN_check_nElement(state, input, gradOutput);
THCUNN_assertSameGPU(state, 2, input, gradInput);
ptrdiff_t size = THCTensor_(nElement)(state, input);
input = THCTensor_(newContiguous)(state, input);
THCTensor_(resizeAs)(state, gradInput, input);
thrust::device_ptr<scalar_t> input_data(THCTensor_(data)(state, input));
thrust::device_ptr<scalar_t> gradInput_data(THCTensor_(data)(state, gradInput));
thrust::transform(input_data, input_data+size, gradInput_data, l1cost_updateGradInput_functor<scalar_t>());
THCTensor_(free)(state, input);
}
#endif