Skip to content

【AMP OP&Test】add fp16 and bf16 test #51286

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Mar 14, 2023
Merged

【AMP OP&Test】add fp16 and bf16 test #51286

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
77da077
add fp16 and bf16 test
zhiboniu Mar 7, 2023
dbf43d1
update
zhiboniu Mar 10, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
25 changes: 13 additions & 12 deletions paddle/phi/kernels/gpu/group_norm_grad_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -51,8 +51,8 @@ __global__ void GroupNormBackwardGetMeanAndVar(const T* x,
if (x_scale != static_cast<T>(0)) x_scale_inv = static_cast<T>(1.0) / x_scale;
AccT d_mean_data = static_cast<AccT>(0);
AccT d_var_data = static_cast<AccT>(0);
T d_scale_data = static_cast<T>(0);
T d_bias_data = static_cast<T>(0);
AccT d_scale_data = static_cast<AccT>(0);
AccT d_bias_data = static_cast<AccT>(0);

for (int imid = threadIdx.x; imid < imsize; imid += blockDim.x) {
AccT val, dval;
Expand All @@ -67,8 +67,8 @@ __global__ void GroupNormBackwardGetMeanAndVar(const T* x,
d_mean_data += dval * static_cast<AccT>(x_scale);

val = val * static_cast<AccT>(x_scale_inv);
d_bias_data += static_cast<T>(dval);
d_scale_data += static_cast<T>(val * dval);
d_bias_data += dval;
d_scale_data += val * dval;
}
CudaAtomicAddWithWarp(&(d_mean[bid * groups + gid]),
static_cast<AccT>(d_mean_data));
Expand All @@ -77,16 +77,16 @@ __global__ void GroupNormBackwardGetMeanAndVar(const T* x,

if (flags & kHasScale) {
#if CUDA_VERSION >= 11070
phi::CudaAtomicAdd(&(d_scale[ccid]), d_scale_data);
phi::CudaAtomicAdd(&(d_scale[ccid]), static_cast<T>(d_scale_data));
#else
CudaAtomicAddWithWarp(&(d_scale[ccid]), d_scale_data);
CudaAtomicAddWithWarp(&(d_scale[ccid]), static_cast<T>(d_scale_data));
#endif
}
if (flags & kHasBias) {
#if CUDA_VERSION >= 11070
phi::CudaAtomicAdd(&(d_bias[ccid]), d_bias_data);
phi::CudaAtomicAdd(&(d_bias[ccid]), static_cast<T>(d_bias_data));
#else
CudaAtomicAddWithWarp(&(d_bias[ccid]), d_bias_data);
CudaAtomicAddWithWarp(&(d_bias[ccid]), static_cast<T>(d_bias_data));
#endif
}
}
Expand Down Expand Up @@ -128,7 +128,7 @@ __global__ void GroupNormBackward(const T* x,
: static_cast<AccT>(1);
AccT x_bias =
(flags & kHasBias) ? static_cast<AccT>(bias[ccid]) : static_cast<AccT>(0);
AccT x_scale_inv = static_cast<T>(0);
AccT x_scale_inv = static_cast<AccT>(0);
if (x_scale != static_cast<AccT>(0))
x_scale_inv = static_cast<AccT>(1.0) / x_scale;

Expand Down Expand Up @@ -220,7 +220,7 @@ __global__ void GetBackwardParamsCUDAKernel(int imsize,
sum1 += static_cast<AccT>(ds[index]) * scale_v;
sum2 += static_cast<AccT>(db[index]) * scale_v;
const AccT scale_c =
scale == nullptr ? static_cast<AccT>(0) : static_cast<T>(scale[c]);
scale == nullptr ? static_cast<AccT>(0) : static_cast<AccT>(scale[c]);
p1[index] = static_cast<AccT>(scale_c) * var_inv;
}

Expand Down Expand Up @@ -402,7 +402,7 @@ void GroupNormGradKernel(const Context& dev_ctx,
p1_data,
p2_data,
p3_data);
GetXGradientCUDAKernel<T>
GetXGradientCUDAKernel<T, AccT>
<<<grid, threads, 0, dev_ctx.stream()>>>(imsize,
C,
group_size,
Expand All @@ -424,7 +424,7 @@ void GroupNormGradKernel(const Context& dev_ctx,

DenseTensor temp_var;
temp_var.Resize(var.dims());
dev_ctx.template Alloc<T>(&temp_var);
dev_ctx.template Alloc<AccT>(&temp_var);
set_zero_AccT(dev_ctx, &temp_var, static_cast<AccT>(0));
auto* temp_var_data = temp_var.data<AccT>();

Expand Down Expand Up @@ -483,4 +483,5 @@ PD_REGISTER_KERNEL(group_norm_grad,
phi::GroupNormGradKernel,
float,
double,
phi::dtype::bfloat16,
phi::dtype::float16) {}
8 changes: 7 additions & 1 deletion paddle/phi/kernels/gpu/group_norm_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,6 +20,11 @@
#include "paddle/phi/kernels/funcs/math_function.h"
#include "paddle/phi/kernels/gpu/group_norm_utils.h"

#include "paddle/phi/common/bfloat16.h"
#include "paddle/phi/common/data_type.h"
#include "paddle/phi/common/float16.h"
#include "paddle/phi/core/device_context.h"

namespace phi {

template <typename T, typename AccT>
Expand Down Expand Up @@ -124,7 +129,7 @@ void GroupNormKernel(const Context& dev_ctx,
DenseTensor* y,
DenseTensor* mean,
DenseTensor* var) {
using AccT = typename kps::details::MPTypeTrait<T>::Type;
using AccT = typename phi::dtype::MPTypeTrait<T>::Type;
const DataLayout data_layout = phi::StringToDataLayout(data_layout_str);
const auto scale_ptr = scale.get_ptr();
const auto bias_ptr = bias.get_ptr();
Expand Down Expand Up @@ -342,4 +347,5 @@ PD_REGISTER_KERNEL(group_norm,
phi::GroupNormKernel,
float,
double,
phi::dtype::bfloat16,
phi::dtype::float16) {}
Loading