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
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
|
#include "THCTensorMathPointwise.cuh"
struct TensorMaxOp {
__device__ __forceinline__ void operator()(float* out, float* in) {
*out = max(*out, *in);
}
__device__ __forceinline__ void operator()(float* out, float* in1, float* in2) {
*out = max(*in1, *in2);
}
};
void THCudaTensor_cmax(THCState *state, THCudaTensor *self, THCudaTensor *src1, THCudaTensor *src2)
{
THAssert(THCudaTensor_checkGPU(state, 3, self, src1, src2));
THArgCheck(THCudaTensor_nElement(state, src1) ==
THCudaTensor_nElement(state, src2), 2, "sizes do not match");
if (self == src1) {
if (!THC_pointwiseApply2(state, self, src2, TensorMaxOp())) {
THArgCheck(false, 2, CUTORCH_DIM_WARNING);
}
} else {
THCudaTensor_resizeAs(state, self, src1);
if (!THC_pointwiseApply3(state, self, src1, src2, TensorMaxOp())) {
THArgCheck(false, 2, CUTORCH_DIM_WARNING);
}
}
}
struct TensorMinOp {
__device__ __forceinline__ void operator()(float* out, float* in) {
*out = min(*out, *in);
}
__device__ __forceinline__ void operator()(float* out, float* in1, float* in2) {
*out = min(*in1, *in2);
}
};
void THCudaTensor_cmin(THCState *state, THCudaTensor *self, THCudaTensor *src1, THCudaTensor *src2)
{
THAssert(THCudaTensor_checkGPU(state, 3, self, src1, src2));
THArgCheck(THCudaTensor_nElement(state, src1) ==
THCudaTensor_nElement(state, src2), 2, "sizes do not match");
if (self == src1) {
if (!THC_pointwiseApply2(state, self, src2, TensorMinOp())) {
THArgCheck(false, 2, CUTORCH_DIM_WARNING);
}
} else {
THCudaTensor_resizeAs(state, self, src1);
if (!THC_pointwiseApply3(state, self, src1, src2, TensorMinOp())) {
THArgCheck(false, 2, CUTORCH_DIM_WARNING);
}
}
}
struct TensorMaxValueOp {
TensorMaxValueOp(float v) : val(v) {}
__device__ __forceinline__ void operator()(float* out) {
*out = max(*out, val);
}
__device__ __forceinline__ void operator()(float* out, float* in) {
*out = max(*in, val);
}
float val;
};
void THCudaTensor_cmaxValue(THCState *state, THCudaTensor *self, THCudaTensor *src, float value)
{
THAssert(THCudaTensor_checkGPU(state, 2, self, src));
if (self == src) {
if (!THC_pointwiseApply1(state, self, TensorMaxValueOp(value))) {
THArgCheck(false, 2, CUTORCH_DIM_WARNING);
}
} else {
THCudaTensor_resizeAs(state, self, src);
if (!THC_pointwiseApply2(state, self, src, TensorMaxValueOp(value))) {
THArgCheck(false, 2, CUTORCH_DIM_WARNING);
}
}
}
struct TensorMinValueOp {
TensorMinValueOp(float v) : val(v) {}
__device__ __forceinline__ void operator()(float* out) {
*out = min(*out, val);
}
__device__ __forceinline__ void operator()(float* out, float* in) {
*out = min(*in, val);
}
float val;
};
void THCudaTensor_cminValue(THCState *state, THCudaTensor *self, THCudaTensor *src, float value)
{
THAssert(THCudaTensor_checkGPU(state, 2, self, src));
if (self == src) {
if (!THC_pointwiseApply1(state, self, TensorMinValueOp(value))) {
THArgCheck(false, 2, CUTORCH_DIM_WARNING);
}
} else {
THCudaTensor_resizeAs(state, self, src);
if (!THC_pointwiseApply2(state, self, src, TensorMinValueOp(value))) {
THArgCheck(false, 2, CUTORCH_DIM_WARNING);
}
}
}
|
