wehub-resource-sync
231 lines
8.0 KiB
Plaintext
231 lines
8.0 KiB
Plaintext
// Copyright (c) 2022 PaddlePaddle Authors. All Rights Reserved.
|
|
//
|
|
// Licensed under the Apache License, Version 2.0 (the "License");
|
|
// you may not use this file except in compliance with the License.
|
|
// You may obtain a copy of the License at
|
|
//
|
|
// http://www.apache.org/licenses/LICENSE-2.0
|
|
//
|
|
// Unless required by applicable law or agreed to in writing, software
|
|
// distributed under the License is distributed on an "AS IS" BASIS,
|
|
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
// See the License for the specific language governing permissions and
|
|
// limitations under the License.
|
|
|
|
#include "paddle/phi/kernels/cross_grad_kernel.h"
|
|
|
|
#include "paddle/phi/backends/gpu/gpu_context.h"
|
|
#include "paddle/phi/backends/gpu/gpu_launch_config.h"
|
|
#include "paddle/phi/core/dense_tensor.h"
|
|
#include "paddle/phi/core/kernel_registry.h"
|
|
#include "paddle/phi/kernels/funcs/complex_functors.h"
|
|
#include "paddle/phi/kernels/funcs/for_range.h"
|
|
#include "paddle/phi/kernels/funcs/index_calculator.h"
|
|
|
|
namespace phi {
|
|
|
|
template <typename T>
|
|
__global__ void CrossGrad(const T* x,
|
|
const T* y,
|
|
const T* out,
|
|
T* out_dx,
|
|
T* out_dy,
|
|
const int64_t stride,
|
|
const int64_t N,
|
|
funcs::IndexCalculator<int64_t> index_calculator) {
|
|
CUDA_KERNEL_LOOP_TYPE(i, N, int64_t) {
|
|
int64_t offset = index_calculator(i);
|
|
|
|
int64_t pos0 = offset + 0 * stride;
|
|
int64_t pos1 = offset + 1 * stride;
|
|
int64_t pos2 = offset + 2 * stride;
|
|
|
|
using MT = typename MPTypeTrait<T>::Type;
|
|
|
|
MT x_pos0_mp = static_cast<MT>(x[pos0]);
|
|
MT x_pos1_mp = static_cast<MT>(x[pos1]);
|
|
MT x_pos2_mp = static_cast<MT>(x[pos2]);
|
|
MT y_pos0_mp = static_cast<MT>(y[pos0]);
|
|
MT y_pos1_mp = static_cast<MT>(y[pos1]);
|
|
MT y_pos2_mp = static_cast<MT>(y[pos2]);
|
|
MT out_pos0_mp = static_cast<MT>(out[pos0]);
|
|
MT out_pos1_mp = static_cast<MT>(out[pos1]);
|
|
MT out_pos2_mp = static_cast<MT>(out[pos2]);
|
|
|
|
out_dx[pos0] =
|
|
static_cast<T>(out_pos2_mp * y_pos1_mp - out_pos1_mp * y_pos2_mp);
|
|
out_dy[pos0] =
|
|
static_cast<T>(out_pos1_mp * x_pos2_mp - out_pos2_mp * x_pos1_mp);
|
|
|
|
out_dx[pos1] =
|
|
static_cast<T>(out_pos0_mp * y_pos2_mp - out_pos2_mp * y_pos0_mp);
|
|
out_dy[pos1] =
|
|
static_cast<T>(out_pos2_mp * x_pos0_mp - out_pos0_mp * x_pos2_mp);
|
|
|
|
out_dx[pos2] =
|
|
static_cast<T>(out_pos1_mp * y_pos0_mp - out_pos0_mp * y_pos1_mp);
|
|
out_dy[pos2] =
|
|
static_cast<T>(out_pos0_mp * x_pos1_mp - out_pos1_mp * x_pos0_mp);
|
|
}
|
|
}
|
|
|
|
template <typename T, typename Context>
|
|
void CrossGradKernel(const Context& dev_ctx,
|
|
const DenseTensor& x,
|
|
const DenseTensor& y,
|
|
const DenseTensor& out_grad,
|
|
int axis,
|
|
DenseTensor* x_grad,
|
|
DenseTensor* y_grad) {
|
|
auto& input_x = x;
|
|
auto& input_y = y;
|
|
auto& input_out_grad = out_grad;
|
|
auto* output_x_grad = x_grad;
|
|
auto* output_y_grad = y_grad;
|
|
int dim = axis;
|
|
|
|
auto input_x_dims = input_x.dims();
|
|
if (dim != DDim::kMaxRank) {
|
|
PADDLE_ENFORCE_EQ(
|
|
dim < input_x_dims.size() && dim >= (0 - input_x_dims.size()),
|
|
true,
|
|
errors::OutOfRange(
|
|
"Attr(dim) is out of range, It's expected "
|
|
"to be in range of [-%d, %d]. But received Attr(dim) = %d.",
|
|
input_x_dims.size(),
|
|
input_x_dims.size() - 1,
|
|
dim));
|
|
if (dim < 0) {
|
|
dim += input_x_dims.size();
|
|
}
|
|
|
|
PADDLE_ENFORCE_EQ(
|
|
input_x_dims[dim] == 3,
|
|
true,
|
|
errors::InvalidArgument(
|
|
"Input(X/Y).dims[dim] must be equal to 3. But received: "
|
|
"Input(X/Y).dims[dim] = [%d].",
|
|
input_x_dims[dim]));
|
|
} else {
|
|
for (auto i = 0; i < input_x_dims.size(); i++) {
|
|
if (input_x_dims[i] == 3) {
|
|
dim = i;
|
|
break;
|
|
}
|
|
}
|
|
PADDLE_ENFORCE_EQ(
|
|
dim == DDim::kMaxRank,
|
|
false,
|
|
errors::InvalidArgument("There must be at least one dimension 'd' "
|
|
"so that Input(X/Y).dims()[d] is equal to 3. "
|
|
"But received: Input(X/Y).dims() == [%s].",
|
|
input_x_dims));
|
|
}
|
|
|
|
std::vector<int64_t> cal_dims;
|
|
std::vector<int64_t> left_strides;
|
|
std::vector<int64_t> full_strides;
|
|
std::vector<int64_t> merged_dims;
|
|
|
|
for (int i = 0; i < dim; i++) {
|
|
if (i == 0) {
|
|
merged_dims.push_back(input_x_dims[i]);
|
|
} else {
|
|
merged_dims[0] *= input_x_dims[i];
|
|
}
|
|
}
|
|
int merge_axis = merged_dims.size();
|
|
merged_dims.push_back(input_x_dims[dim]);
|
|
for (int i = dim + 1; i < input_x_dims.size(); i++) {
|
|
if (i == dim + 1) {
|
|
merged_dims.push_back(input_x_dims[i]);
|
|
} else {
|
|
merged_dims[merge_axis + 1] *= input_x_dims[i];
|
|
}
|
|
}
|
|
|
|
int64_t full_dim = 1;
|
|
for (int i = 0; i < merged_dims.size(); i++) {
|
|
full_strides.insert(full_strides.begin(), full_dim);
|
|
full_dim *= merged_dims[merged_dims.size() - i - 1];
|
|
if (i == merge_axis) {
|
|
continue;
|
|
}
|
|
cal_dims.push_back(i);
|
|
}
|
|
int64_t left_dim = 1;
|
|
for (int i = merged_dims.size() - 1; i >= 0; i--) {
|
|
if (i == merge_axis) {
|
|
continue;
|
|
}
|
|
left_strides.insert(left_strides.begin(), left_dim);
|
|
left_dim *= merged_dims[i];
|
|
}
|
|
|
|
const auto* input_x_data = input_x.data<T>();
|
|
const auto* input_y_data = input_y.data<T>();
|
|
int64_t numel = x.numel();
|
|
const auto* input_out_grad_data = input_out_grad.data<T>();
|
|
auto* output_x_grad_data = dev_ctx.template Alloc<T>(x_grad);
|
|
auto* output_y_grad_data = dev_ctx.template Alloc<T>(y_grad);
|
|
if (numel == 0) {
|
|
return;
|
|
}
|
|
backends::gpu::GpuLaunchConfig config =
|
|
backends::gpu::GetGpuLaunchConfig1D(dev_ctx, numel / 3);
|
|
auto index_calculator = funcs::IndexCalculator<int64_t>(
|
|
merged_dims.size() - 1, cal_dims, left_strides, full_strides);
|
|
if (IsComplexType(x.dtype())) {
|
|
DenseTensor x_conj, y_conj;
|
|
DenseTensorMeta meta_xy(x.dtype(), x.dims());
|
|
x_conj.set_meta(meta_xy);
|
|
y_conj.set_meta(meta_xy);
|
|
|
|
auto* input_x_conj_data = dev_ctx.template Alloc<T>(&x_conj);
|
|
auto* input_y_conj_data = dev_ctx.template Alloc<T>(&y_conj);
|
|
|
|
funcs::ForRange<Context> for_range(dev_ctx, numel);
|
|
funcs::ConjFunctor<T> functor_x(input_x_data, numel, input_x_conj_data);
|
|
funcs::ConjFunctor<T> functor_y(input_y_data, numel, input_y_conj_data);
|
|
for_range(functor_x);
|
|
for_range(functor_y);
|
|
CrossGrad<<<config.block_per_grid,
|
|
config.thread_per_block,
|
|
0,
|
|
dev_ctx.stream()>>>(input_x_conj_data,
|
|
input_y_conj_data,
|
|
input_out_grad_data,
|
|
output_x_grad_data,
|
|
output_y_grad_data,
|
|
full_strides[merge_axis],
|
|
static_cast<int64_t>(numel / 3),
|
|
index_calculator);
|
|
} else {
|
|
CrossGrad<<<config.block_per_grid,
|
|
config.thread_per_block,
|
|
0,
|
|
dev_ctx.stream()>>>(input_x_data,
|
|
input_y_data,
|
|
input_out_grad_data,
|
|
output_x_grad_data,
|
|
output_y_grad_data,
|
|
full_strides[merge_axis],
|
|
static_cast<int64_t>(numel / 3),
|
|
index_calculator);
|
|
}
|
|
}
|
|
} // namespace phi
|
|
|
|
PD_REGISTER_KERNEL(cross_grad,
|
|
GPU,
|
|
ALL_LAYOUT,
|
|
phi::CrossGradKernel,
|
|
phi::float16,
|
|
phi::bfloat16,
|
|
float,
|
|
double,
|
|
int,
|
|
int64_t,
|
|
phi::complex64,
|
|
phi::complex128) {}
|