|
|
|
@ -21,9 +21,9 @@ namespace math {
|
|
|
|
|
|
|
|
|
|
template <typename T>
|
|
|
|
|
__global__ void KernelMaxOut(const int nthreads, const T* input_data,
|
|
|
|
|
const int channels,
|
|
|
|
|
const int input_height, const int input_width,
|
|
|
|
|
int groups, T* output_data ) {
|
|
|
|
|
const int channels, const int input_height,
|
|
|
|
|
const int input_width, int groups,
|
|
|
|
|
T* output_data) {
|
|
|
|
|
const int size = input_height * input_width * channels / groups;
|
|
|
|
|
const int feat_len = input_height * input_width;
|
|
|
|
|
int index = blockIdx.x * blockDim.x + threadIdx.x;
|
|
|
|
@ -34,7 +34,7 @@ __global__ void KernelMaxOut(const int nthreads, const T* input_data,
|
|
|
|
|
int channel_idx = batch_offset / feat_len;
|
|
|
|
|
int feat_idx = batch_offset % feat_len;
|
|
|
|
|
int data_idx =
|
|
|
|
|
(batch_idx * size + channel_idx * feat_len) * groups + feat_idx;
|
|
|
|
|
(batch_idx * size + channel_idx * feat_len) * groups + feat_idx;
|
|
|
|
|
T ele = static_cast<T>(-FLT_MAX);
|
|
|
|
|
for (int g = 0; g < groups; ++g) {
|
|
|
|
|
T x = input_data[data_idx + g * feat_len];
|
|
|
|
@ -44,34 +44,35 @@ __global__ void KernelMaxOut(const int nthreads, const T* input_data,
|
|
|
|
|
}
|
|
|
|
|
}
|
|
|
|
|
template <typename T>
|
|
|
|
|
__global__ void KernelMaxoutGrad(
|
|
|
|
|
const int nthreads, const T* input_data, const T* output_data,
|
|
|
|
|
const T* output_grad, T* input_grad, const int channels,
|
|
|
|
|
const int input_height, const int input_width, int groups) {
|
|
|
|
|
const int size = input_height * input_width * channels / groups;
|
|
|
|
|
const int feat_len = input_height * input_width;
|
|
|
|
|
int index = blockIdx.x * blockDim.x + threadIdx.x;
|
|
|
|
|
int offset = blockDim.x * gridDim.x;
|
|
|
|
|
for (int i = index; i < nthreads; i += offset) {
|
|
|
|
|
int batch_idx = i / size;
|
|
|
|
|
int batch_offset = i % size;
|
|
|
|
|
int channel_idx = batch_offset / feat_len;
|
|
|
|
|
int feat_idx = batch_offset % feat_len;
|
|
|
|
|
int data_idx =
|
|
|
|
|
__global__ void KernelMaxoutGrad(const int nthreads, const T* input_data,
|
|
|
|
|
const T* output_data, const T* output_grad,
|
|
|
|
|
T* input_grad, const int channels,
|
|
|
|
|
const int input_height, const int input_width,
|
|
|
|
|
int groups) {
|
|
|
|
|
const int size = input_height * input_width * channels / groups;
|
|
|
|
|
const int feat_len = input_height * input_width;
|
|
|
|
|
int index = blockIdx.x * blockDim.x + threadIdx.x;
|
|
|
|
|
int offset = blockDim.x * gridDim.x;
|
|
|
|
|
for (int i = index; i < nthreads; i += offset) {
|
|
|
|
|
int batch_idx = i / size;
|
|
|
|
|
int batch_offset = i % size;
|
|
|
|
|
int channel_idx = batch_offset / feat_len;
|
|
|
|
|
int feat_idx = batch_offset % feat_len;
|
|
|
|
|
int data_idx =
|
|
|
|
|
(batch_idx * size + channel_idx * feat_len) * groups + feat_idx;
|
|
|
|
|
int max_index = -1;
|
|
|
|
|
bool continue_match = true;
|
|
|
|
|
for (int g = 0; g < groups && continue_match; ++g) {
|
|
|
|
|
if (input_data[data_idx + g * feat_len] == output_data[i]) {
|
|
|
|
|
max_index = data_idx + g * feat_len;
|
|
|
|
|
continue_match = false;
|
|
|
|
|
break;
|
|
|
|
|
}
|
|
|
|
|
}
|
|
|
|
|
if (max_index != -1) {
|
|
|
|
|
input_grad[max_index] += output_grad[index];
|
|
|
|
|
int max_index = -1;
|
|
|
|
|
bool continue_match = true;
|
|
|
|
|
for (int g = 0; g < groups && continue_match; ++g) {
|
|
|
|
|
if (input_data[data_idx + g * feat_len] == output_data[i]) {
|
|
|
|
|
max_index = data_idx + g * feat_len;
|
|
|
|
|
continue_match = false;
|
|
|
|
|
break;
|
|
|
|
|
}
|
|
|
|
|
}
|
|
|
|
|
if (max_index != -1) {
|
|
|
|
|
input_grad[max_index] += output_grad[index];
|
|
|
|
|
}
|
|
|
|
|
}
|
|
|
|
|
}
|
|
|
|
|
/*
|
|
|
|
|
* All tensors are in NCHW format.
|
|
|
|
@ -80,7 +81,7 @@ template <typename T>
|
|
|
|
|
class MaxOutFunctor<platform::GPUPlace, T> {
|
|
|
|
|
public:
|
|
|
|
|
void operator()(const platform::DeviceContext& context,
|
|
|
|
|
const framework::Tensor& input, framework::Tensor * output,
|
|
|
|
|
const framework::Tensor& input, framework::Tensor* output,
|
|
|
|
|
int groups) {
|
|
|
|
|
const int batch_size = input.dims()[0];
|
|
|
|
|
const int input_channels = input.dims()[1];
|
|
|
|
@ -92,7 +93,7 @@ class MaxOutFunctor<platform::GPUPlace, T> {
|
|
|
|
|
|
|
|
|
|
const T* input_data = input.data<T>();
|
|
|
|
|
T* output_data = output->mutable_data<T>(context.GetPlace());
|
|
|
|
|
int nthreads = output->numel();
|
|
|
|
|
int nthreads = output->numel();
|
|
|
|
|
int blocks = (nthreads + 1024 - 1) / 1024;
|
|
|
|
|
dim3 threads(1024, 1);
|
|
|
|
|
dim3 grid(blocks, 1);
|
|
|
|
@ -101,8 +102,7 @@ class MaxOutFunctor<platform::GPUPlace, T> {
|
|
|
|
|
T><<<grid, threads, 0,
|
|
|
|
|
reinterpret_cast<const platform::CUDADeviceContext&>(context)
|
|
|
|
|
.stream()>>>(nthreads, input_data, input_channels,
|
|
|
|
|
input_height, input_width, groups,
|
|
|
|
|
output_data);
|
|
|
|
|
input_height, input_width, groups, output_data);
|
|
|
|
|
}
|
|
|
|
|
};
|
|
|
|
|
/*
|
|
|
|
@ -112,11 +112,9 @@ template <typename T>
|
|
|
|
|
class MaxOutGradFunctor<platform::GPUPlace, T> {
|
|
|
|
|
public:
|
|
|
|
|
void operator()(const platform::DeviceContext& context,
|
|
|
|
|
const framework::Tensor& input,
|
|
|
|
|
framework::Tensor * input_grad,
|
|
|
|
|
const framework::Tensor& input, framework::Tensor* input_grad,
|
|
|
|
|
const framework::Tensor& output,
|
|
|
|
|
const framework::Tensor& output_grad,
|
|
|
|
|
int groups) {
|
|
|
|
|
const framework::Tensor& output_grad, int groups) {
|
|
|
|
|
const int batch_size = input.dims()[0];
|
|
|
|
|
const int input_channels = input.dims()[1];
|
|
|
|
|
const int input_height = input.dims()[2];
|
|
|
|
@ -129,7 +127,7 @@ class MaxOutGradFunctor<platform::GPUPlace, T> {
|
|
|
|
|
const T* output_data = output.data<T>();
|
|
|
|
|
const T* output_grad_data = output_grad.data<T>();
|
|
|
|
|
T* input_grad_data = input_grad->mutable_data<T>(context.GetPlace());
|
|
|
|
|
int nthreads = output.numel();
|
|
|
|
|
int nthreads = output.numel();
|
|
|
|
|
int blocks = (nthreads + 1024 - 1) / 1024;
|
|
|
|
|
dim3 threads(1024, 1);
|
|
|
|
|
dim3 grid(blocks, 1);
|
|
|
|
@ -137,9 +135,9 @@ class MaxOutGradFunctor<platform::GPUPlace, T> {
|
|
|
|
|
KernelMaxoutGrad<
|
|
|
|
|
T><<<grid, threads, 0,
|
|
|
|
|
reinterpret_cast<const platform::CUDADeviceContext&>(context)
|
|
|
|
|
.stream()>>>(
|
|
|
|
|
nthreads, input_data, output_data, output_grad_data, input_grad_data,
|
|
|
|
|
input_channels, input_height, input_width, groups);
|
|
|
|
|
.stream()>>>(nthreads, input_data, output_data,
|
|
|
|
|
output_grad_data, input_grad_data, input_channels,
|
|
|
|
|
input_height, input_width, groups);
|
|
|
|
|
}
|
|
|
|
|
};
|
|
|
|
|
|
|
|
|
|