Skip to content

Commit d5172d9

Browse files
YqGe585tianshuo78520azrr1999duqimeng
authored
Fix top_p_sampling kernel (#2231)
* update paddle * fix metax error * fix fusion error * fix metax error * fix cmake * fix cmake * update paddle to 1013 * fix patch * fix patch * update paddle * fix top_p_sampling_kernel * Update Paddle submodule to latest develop * fix scale * fix * fix datatype * restore cmake * fix patch * [Metax] Fix DNN-related bugs (#194) [Metax] Fix DNN-related bugs (#194) * fix patch * fix compile error * fix scaled_dot_product_attention * disable test_tensor_cuda_static --------- Co-authored-by: tianshuo78520a <tianshuo78520a@users.noreply.github.com> Co-authored-by: zrr1999 <2742392377@qq.com> Co-authored-by: duqimeng <77875733+duqimeng@users.noreply.github.com>
1 parent 7deecbf commit d5172d9

33 files changed

Lines changed: 306 additions & 347 deletions

‎Paddle‎

Submodule Paddle updated 655 files

‎backends/gcu/kernels/interpolate_kernels.cc‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -28,7 +28,7 @@ void InterpolateKernel(
2828
int out_d,
2929
int out_h,
3030
int out_w,
31-
const std::vector<float>& scale,
31+
const std::vector<double>& scale,
3232
const std::string& interp_method,
3333
bool align_corners,
3434
int align_mode,
@@ -47,7 +47,7 @@ void InterpolateKernel(
4747

4848
float scale_h = -1;
4949
float scale_w = -1;
50-
std::vector<float> new_scale(scale);
50+
std::vector<double> new_scale(scale);
5151
// Priority: size_tensor > out_size > scale_tensor > scale > out_h & out_w
5252
if (size_tensor && size_tensor->size() > 0) {
5353
auto tensors = size_tensor.get();
@@ -253,7 +253,7 @@ void InterpolateGradKernel(
253253
int out_d,
254254
int out_h,
255255
int out_w,
256-
const std::vector<float>& scale,
256+
const std::vector<double>& scale,
257257
const std::string& interp_method,
258258
bool align_corners,
259259
int align_mode,
@@ -326,7 +326,7 @@ void BilinearInterpKernel(
326326
int out_d,
327327
int out_h,
328328
int out_w,
329-
const std::vector<float>& scale,
329+
const std::vector<double>& scale,
330330
const std::string& interp_method,
331331
bool align_corners,
332332
int align_mode,
@@ -361,7 +361,7 @@ void BilinearInterpGradKernel(
361361
int out_d,
362362
int out_h,
363363
int out_w,
364-
const std::vector<float>& scale,
364+
const std::vector<double>& scale,
365365
const std::string& interp_method,
366366
bool align_corners,
367367
int align_mode,
@@ -400,7 +400,7 @@ void NearestInterpKernel(
400400
int out_d,
401401
int out_h,
402402
int out_w,
403-
const std::vector<float>& scale,
403+
const std::vector<double>& scale,
404404
const std::string& interp_method,
405405
bool align_corners,
406406
int align_mode,
@@ -435,7 +435,7 @@ void NearestInterpGradKernel(
435435
int out_d,
436436
int out_h,
437437
int out_w,
438-
const std::vector<float>& scale,
438+
const std::vector<double>& scale,
439439
const std::string& interp_method,
440440
bool align_corners,
441441
int align_mode,

‎backends/iluvatar_gpu/build_inc.sh‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -150,7 +150,7 @@ fi
150150

151151
# Compile
152152
echo "Starting compilation..."
153-
ninja -k 0 -j$(nproc) 2>&1 | tee -a compile.log
153+
ninja -j$(nproc) 2>&1
154154
FAILED_LOG="failed_files.log"
155155
grep -E "FAILED: " compile.log | tee ${FAILED_LOG}
156156
echo "Failed files are listed in ${FAILED_LOG}"

‎backends/iluvatar_gpu/build_paddle.sh‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -72,7 +72,7 @@ cmake -G Ninja -DPY_VERSION=${PYTHON_VERSION} -DWITH_COREX=ON -DPADDLE_SOURCE_DI
7272
-DCMAKE_CUDA_FLAGS='-Xclang -fcuda-allow-variadic-functions -mllvm --skip-double' \
7373
-DCMAKE_C_FLAGS="-pthread" \
7474
-DWITH_ARM=OFF -DWITH_DGC=OFF .. || { echo "Error: CMake configuration failed!"; exit 1; }
75-
ninja -k 0 -j$(nproc) || { echo "Error: Paddle-iluvatar-gpu build failed!"; exit 1; }
75+
ninja -j$(nproc) || { echo "Error: Paddle-iluvatar-gpu build failed!"; exit 1; }
7676
popd
7777

7878
if [[ ! -d "build_pip" ]]; then

‎backends/iluvatar_gpu/kernels/cuda_kernels/cross_entropy_kernel.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -751,7 +751,7 @@ static void SoftmaxWithCrossEntropySoftLabel(const GPUContext& dev_ctx,
751751
} else {
752752
ScopedTensorDescriptor desc;
753753
std::vector<int> tensor_dims = {N, dim, D, 1};
754-
GPUDNNDataLayout layout = GPUDNNDataLayout::kNCHW;
754+
DataLayout layout = DataLayout::kNCHW;
755755
cudnnTensorDescriptor_t descp = desc.descriptor<T>(layout, tensor_dims);
756756

757757
auto handle = GetDnnHandle(dev_ctx.stream(), dev_ctx.GetPlace());
@@ -1163,7 +1163,7 @@ static void SoftmaxWithCrossEntropyHardLabel(const GPUContext& dev_ctx,
11631163
} else {
11641164
ScopedTensorDescriptor desc;
11651165
std::vector<int> tensor_dims = {N, dim, D, 1};
1166-
GPUDNNDataLayout layout = GPUDNNDataLayout::kNCHW;
1166+
DataLayout layout = DataLayout::kNCHW;
11671167
cudnnTensorDescriptor_t descp = desc.descriptor<T>(layout, tensor_dims);
11681168
auto handle = GetDnnHandle(dev_ctx.stream(), dev_ctx.GetPlace());
11691169

‎backends/iluvatar_gpu/kernels/ernie_core/top_p_sampling_kernel.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1054,7 +1054,7 @@ void TopPSamplingKernel(const Context& dev_ctx,
10541054
const DenseTensor& ps,
10551055
const paddle::optional<DenseTensor>& threshold,
10561056
const paddle::optional<DenseTensor>& topp_seed,
1057-
int seed,
1057+
int64_t seed,
10581058
int k,
10591059
const std::string& mode,
10601060
DenseTensor* out,

‎backends/iluvatar_gpu/kernels/gpudnn/conv_grad_kernel.cu‎

Lines changed: 31 additions & 38 deletions
Original file line numberDiff line numberDiff line change
@@ -53,8 +53,8 @@ void ConvCudnnGradKernelImplV7(
5353
const std::vector<int>& strides,
5454
const std::vector<int>& padding_common,
5555
const std::vector<int>& dilations,
56-
phi::backends::gpu::DataLayout compute_format,
57-
phi::backends::gpu::DataLayout layout,
56+
DataLayout compute_format,
57+
DataLayout layout,
5858
bool use_addto,
5959
bool exhaustive_search,
6060
bool deterministic,
@@ -98,31 +98,31 @@ void ConvCudnnGradKernelImplV7(
9898

9999
int i_n, i_c, i_d, i_h, i_w;
100100
int o_n, o_c, o_d, o_h, o_w;
101-
if (compute_format == phi::backends::gpu::DataLayout::kNHWC) {
101+
if (compute_format == DataLayout::NHWC) {
102102
GetNCDHW(transformed_input->dims(),
103-
phi::backends::gpu::DataLayout::kNHWC,
103+
DataLayout::NHWC,
104104
&i_n,
105105
&i_c,
106106
&i_d,
107107
&i_h,
108108
&i_w);
109109
GetNCDHW(transformed_output_grad_channel->dims(),
110-
phi::backends::gpu::DataLayout::kNHWC,
110+
DataLayout::NHWC,
111111
&o_n,
112112
&o_c,
113113
&o_d,
114114
&o_h,
115115
&o_w);
116116
} else {
117117
GetNCDHW(transformed_input->dims(),
118-
phi::backends::gpu::DataLayout::kNCHW,
118+
DataLayout::NCHW,
119119
&i_n,
120120
&i_c,
121121
&i_d,
122122
&i_h,
123123
&i_w);
124124
GetNCDHW(transformed_output_grad_channel->dims(),
125-
phi::backends::gpu::DataLayout::kNCHW,
125+
DataLayout::NCHW,
126126
&o_n,
127127
&o_c,
128128
&o_d,
@@ -349,7 +349,7 @@ void ConvCudnnGradKernelImplV8(
349349
const std::vector<int>& strides,
350350
const std::vector<int>& padding_common,
351351
const std::vector<int>& dilations,
352-
phi::backends::gpu::DataLayout layout,
352+
DataLayout layout,
353353
bool use_addto,
354354
bool exhaustive_search,
355355
bool deterministic,
@@ -469,7 +469,7 @@ void ConvCudnnGradKernel(const Context& dev_ctx,
469469

470470
#ifdef PADDLE_WITH_HIP
471471
// HIP MIOPEN ONLY SUPPORT NCHW format
472-
auto compute_format = phi::backends::gpu::DataLayout::kNCHW;
472+
auto compute_format = DataLayout::NCHW;
473473
#else
474474
#if CUDNN_VERSION_MIN(8, 1, 0)
475475
const bool compute_in_nhwc =
@@ -479,14 +479,12 @@ void ConvCudnnGradKernel(const Context& dev_ctx,
479479
const bool compute_in_nhwc =
480480
dtype == CUDNN_DATA_HALF && IsVoltaOrLater(dev_ctx);
481481
#endif
482-
auto compute_format = compute_in_nhwc && channel_last
483-
? phi::backends::gpu::DataLayout::kNHWC
484-
: phi::backends::gpu::DataLayout::kNCHW;
482+
auto compute_format =
483+
compute_in_nhwc && channel_last ? DataLayout::NHWC : DataLayout::NCHW;
485484
#endif
486485
VLOG(3) << "Compute ConvGradOp with cuDNN:"
487486
<< " data_format=" << data_format << " compute_format="
488-
<< (compute_format == phi::backends::gpu::DataLayout::kNHWC ? "NHWC"
489-
: "NCHW");
487+
<< (compute_format == DataLayout::NHWC ? "NHWC" : "NCHW");
490488

491489
// transform Tensor
492490
DenseTensor transformed_input_channel(input.type());
@@ -495,7 +493,7 @@ void ConvCudnnGradKernel(const Context& dev_ctx,
495493
DenseTensor transformed_filter_channel(filter.type());
496494
DenseTensor transformed_filter_grad_channel(filter.type());
497495

498-
if (channel_last && compute_format == phi::backends::gpu::DataLayout::kNCHW) {
496+
if (channel_last && compute_format == DataLayout::NCHW) {
499497
VLOG(3) << "Transform input, output_grad, input_grad and tensor from "
500498
"NHWC to NCHW.";
501499
ResizeToChannelFirst<Context, T>(
@@ -526,7 +524,7 @@ void ConvCudnnGradKernel(const Context& dev_ctx,
526524
}
527525
}
528526

529-
if (compute_format == phi::backends::gpu::DataLayout::kNHWC) {
527+
if (compute_format == DataLayout::NHWC) {
530528
VLOG(3) << "Transform filter and filter_grad tensor from NCHW to NHWC.";
531529
ResizeToChannelLast<Context, T>(
532530
dev_ctx, &filter, &transformed_filter_channel);
@@ -549,7 +547,7 @@ void ConvCudnnGradKernel(const Context& dev_ctx,
549547
auto filter_dims = transformed_filter_channel.dims();
550548
DDim in_data_dims;
551549
DDim filter_data_dims;
552-
if (compute_format == phi::backends::gpu::DataLayout::kNCHW) {
550+
if (compute_format == DataLayout::NCHW) {
553551
in_data_dims = slice_ddim(in_dims, 2, in_dims.size());
554552
filter_data_dims = slice_ddim(filter_dims, 2, filter_dims.size());
555553
} else {
@@ -574,7 +572,7 @@ void ConvCudnnGradKernel(const Context& dev_ctx,
574572
std::vector<int> padding_diff(data_dim);
575573
std::vector<int> new_input_shape_vec(data_dim + 2);
576574
new_input_shape_vec[0] = transformed_input_channel.dims()[0];
577-
if (compute_format == phi::backends::gpu::DataLayout::kNCHW) {
575+
if (compute_format == DataLayout::NCHW) {
578576
new_input_shape_vec[1] = transformed_input_channel.dims()[1];
579577
} else {
580578
new_input_shape_vec[data_dim + 1] =
@@ -584,14 +582,14 @@ void ConvCudnnGradKernel(const Context& dev_ctx,
584582
for (size_t i = 0; i < data_dim; ++i) {
585583
padding_diff[i] = std::abs(paddings[2 * i] - paddings[2 * i + 1]);
586584
padding_common[i] = std::min(paddings[2 * i], paddings[2 * i + 1]);
587-
if (compute_format == phi::backends::gpu::DataLayout::kNCHW) {
585+
if (compute_format == DataLayout::NCHW) {
588586
new_input_shape_vec[i + 2] =
589587
transformed_input_channel.dims()[i + 2] + padding_diff[i];
590588
} else {
591589
new_input_shape_vec[i + 1] =
592590
transformed_input_channel.dims()[i + 1] + padding_diff[i];
593591
}
594-
if (compute_format == phi::backends::gpu::DataLayout::kNCHW) {
592+
if (compute_format == DataLayout::NCHW) {
595593
input_pad[2 * i + 4] = paddings[2 * i] - padding_common[i];
596594
input_pad[2 * i + 4 + 1] = paddings[2 * i + 1] - padding_common[i];
597595
} else {
@@ -645,14 +643,11 @@ void ConvCudnnGradKernel(const Context& dev_ctx,
645643
}
646644
}
647645
}
648-
phi::backends::gpu::DataLayout layout =
649-
compute_format == phi::backends::gpu::DataLayout::kNHWC
650-
? phi::backends::gpu::DataLayout::kNHWC
651-
: phi::backends::gpu::DataLayout::kNCHW;
646+
DataLayout layout =
647+
compute_format == DataLayout::NHWC ? DataLayout::NHWC : DataLayout::NCHW;
652648
if (transformed_input.dims().size() == 5) {
653-
layout = compute_format == phi::backends::gpu::DataLayout::kNHWC
654-
? phi::backends::gpu::DataLayout::kNDHWC
655-
: phi::backends::gpu::DataLayout::kNCDHW;
649+
layout = compute_format == DataLayout::NHWC ? DataLayout::NDHWC
650+
: DataLayout::NCDHW;
656651
}
657652
CUDNN_ENFORCE_TENSOR_SIZE_SUPPORTED(transformed_input);
658653
CUDNN_ENFORCE_TENSOR_SIZE_SUPPORTED(transformed_filter_channel);
@@ -740,15 +735,14 @@ void ConvCudnnGradKernel(const Context& dev_ctx,
740735
}
741736
}
742737

743-
if (channel_last &&
744-
compute_format == phi::backends::gpu::DataLayout::kNCHW) {
738+
if (channel_last && compute_format == DataLayout::NCHW) {
745739
TransToChannelLast<Context, T>(
746740
dev_ctx, &transformed_input_grad_channel, input_grad);
747741
}
748742
}
749743

750744
if (filter_grad) {
751-
if (compute_format == phi::backends::gpu::DataLayout::kNHWC) {
745+
if (compute_format == DataLayout::NHWC) {
752746
TransToChannelFirst<Context, T>(
753747
dev_ctx, &transformed_filter_grad_channel, filter_grad);
754748
}
@@ -1011,8 +1005,7 @@ void ConvCudnnGradGradKernel(
10111005
auto dtype = phi::backends::gpu::CudnnDataType<T>::type;
10121006

10131007
auto handle = GetDnnHandle(dev_ctx.stream(), dev_ctx.GetPlace());
1014-
auto layout = phi::backends::gpu::GetCudnnTensorFormat(
1015-
phi::backends::gpu::DataLayout::kNCHW);
1008+
auto layout = phi::backends::gpu::GetCudnnTensorFormat(DataLayout::NCHW);
10161009

10171010
ConvArgs args1{handle,
10181011
&transformed_ddX,
@@ -1023,7 +1016,7 @@ void ConvCudnnGradGradKernel(
10231016
dilations,
10241017
dtype,
10251018
groups,
1026-
phi::backends::gpu::DataLayout::kNCHW};
1019+
DataLayout::NCHW};
10271020
ConvArgs args2{handle,
10281021
&transformed_X,
10291022
ddW,
@@ -1033,7 +1026,7 @@ void ConvCudnnGradGradKernel(
10331026
dilations,
10341027
dtype,
10351028
groups,
1036-
phi::backends::gpu::DataLayout::kNCHW};
1029+
DataLayout::NCHW};
10371030
ConvArgs args3{handle,
10381031
&transformed_ddX,
10391032
dW,
@@ -1043,7 +1036,7 @@ void ConvCudnnGradGradKernel(
10431036
dilations,
10441037
dtype,
10451038
groups,
1046-
phi::backends::gpu::DataLayout::kNCHW};
1039+
DataLayout::NCHW};
10471040
ConvArgs args4{handle,
10481041
&transformed_dX,
10491042
ddW,
@@ -1053,7 +1046,7 @@ void ConvCudnnGradGradKernel(
10531046
dilations,
10541047
dtype,
10551048
groups,
1056-
phi::backends::gpu::DataLayout::kNCHW};
1049+
DataLayout::NCHW};
10571050

10581051
#ifdef PADDLE_WITH_HIP
10591052
SearchResult<miopenConvFwdAlgorithm_t> fwd_result1;
@@ -1179,11 +1172,11 @@ void ConvCudnnGradGradKernel(
11791172

11801173
int i_n, i_c, i_d, i_h, i_w;
11811174
GetNCDHW(
1182-
transformed_X.dims(), DataLayout::kNCHW, &i_n, &i_c, &i_d, &i_h, &i_w);
1175+
transformed_X.dims(), DataLayout::NCHW, &i_n, &i_c, &i_d, &i_h, &i_w);
11831176

11841177
int o_n, o_c, o_d, o_h, o_w;
11851178
GetNCDHW(transformed_dO_channel.dims(),
1186-
DataLayout::kNCHW,
1179+
DataLayout::NCHW,
11871180
&o_n,
11881181
&o_c,
11891182
&o_d,

0 commit comments

Comments
 (0)