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
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
|
#ifndef THC_GENERIC_FILE
#define THC_GENERIC_FILE "generic/BCECriterion.cu"
#else
void THNN_(BCECriterion_updateOutput)(
THCState *state,
THCTensor *input,
THCTensor *target,
THCTensor *output,
bool sizeAverage,
THCTensor *weights)
{
THCUNN_assertSameGPU_generic(state, 3, input, target, weights);
long size = THCTensor_(nElement)(state, input);
input = THCTensor_(newContiguous)(state, input);
target = THCTensor_(newContiguous)(state, target);
thrust::device_ptr<real> input_data(THCTensor_(data)(state, input));
thrust::device_ptr<real> target_data(THCTensor_(data)(state, target));
accreal sum;
if (weights) {
weights = THCTensor_(newContiguous)(state, weights);
thrust::device_ptr<real> weights_data(THCTensor_(data)(state, weights));
sum = thrust::transform_reduce(
thrust::make_zip_iterator(thrust::make_tuple(input_data, target_data, weights_data)),
thrust::make_zip_iterator(thrust::make_tuple(input_data+size, target_data+size, weights_data+size)),
bce_functor_weights<real, accreal>(),
(accreal) 0,
thrust::plus<accreal>()
);
THCTensor_(free)(state, weights);
} else {
sum = thrust::transform_reduce(
thrust::make_zip_iterator(thrust::make_tuple(input_data, target_data)),
thrust::make_zip_iterator(thrust::make_tuple(input_data+size, target_data+size)),
bce_functor<real, accreal>(),
(accreal) 0,
thrust::plus<accreal>()
);
}
if (sizeAverage)
sum /= size;
THCTensor_(free)(state, input);
THCTensor_(free)(state, target);
THCTensor_(set1d)(state, output, 0, ScalarConvert<accreal, real>::to(sum));
}
void THNN_(BCECriterion_updateGradInput)(
THCState *state,
THCTensor *input,
THCTensor *target,
THCTensor *gradInput,
bool sizeAverage,
THCTensor *weights)
{
THCUNN_assertSameGPU_generic(state, 4, input, target, gradInput, weights);
long size = THCTensor_(nElement)(state, input);
real norm = ScalarConvert<accreal, real>::to(sizeAverage ? accreal(1)/size : accreal(1));
input = THCTensor_(newContiguous)(state, input);
target = THCTensor_(newContiguous)(state, target);
THCTensor_(resizeAs)(state, gradInput, input);
thrust::device_ptr<real> input_data(THCTensor_(data)(state, input));
thrust::device_ptr<real> target_data(THCTensor_(data)(state, target));
thrust::device_ptr<real> gradInput_data(THCTensor_(data)(state, gradInput));
if (weights) {
weights = THCTensor_(newContiguous)(state, weights);
thrust::device_ptr<real> weights_data(THCTensor_(data)(state, weights));
thrust::transform(
thrust::make_zip_iterator(thrust::make_tuple(input_data, target_data, weights_data)),
thrust::make_zip_iterator(thrust::make_tuple(input_data+size, target_data+size, weights_data+size)),
gradInput_data,
bce_updateGradInput_functor_weights<real, accreal>(norm)
);
THCTensor_(free)(state, weights);
} else {
thrust::transform(
thrust::make_zip_iterator(thrust::make_tuple(input_data, target_data)),
thrust::make_zip_iterator(thrust::make_tuple(input_data+size, target_data+size)),
gradInput_data,
bce_updateGradInput_functor<real, accreal>(norm)
);
}
THCTensor_(free)(state, input);
THCTensor_(free)(state, target);
}
#endif
|
