wehub-resource-sync
107 lines
3.8 KiB
Plaintext
107 lines
3.8 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/shard_index_kernel.h"
|
|
|
|
#include "paddle/phi/backends/gpu/gpu_context.h"
|
|
#include "paddle/phi/backends/gpu/gpu_primitives.h"
|
|
#include "paddle/phi/core/kernel_registry.h"
|
|
|
|
namespace phi {
|
|
|
|
template <typename T>
|
|
__global__ void ShardIndexInner(const T* in_data,
|
|
T* out_data,
|
|
const int64_t numel,
|
|
const int index_num,
|
|
const int nshards,
|
|
const int shard_id,
|
|
const int ignore_value) {
|
|
int shard_size = (index_num + nshards - 1) / nshards;
|
|
int64_t idx = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
|
|
if (idx < numel) {
|
|
PADDLE_ENFORCE(in_data[idx] >= 0,
|
|
"The input_index for Op(shard_index) must be "
|
|
"greater or equal to 0, but the value given is %d.",
|
|
in_data[idx]);
|
|
PADDLE_ENFORCE(in_data[idx] < index_num,
|
|
"The input_index for Op(shard_index) must be less "
|
|
"than index_num (%d), but the value given is %d.",
|
|
index_num,
|
|
in_data[idx]);
|
|
if (in_data[idx] / shard_size == shard_id) {
|
|
out_data[idx] = in_data[idx] % shard_size;
|
|
} else {
|
|
out_data[idx] = ignore_value;
|
|
}
|
|
}
|
|
}
|
|
|
|
template <typename T, typename Context>
|
|
void ShardIndexKernel(const Context& dev_ctx,
|
|
const DenseTensor& in,
|
|
int index_num,
|
|
int nshards,
|
|
int shard_id,
|
|
int ignore_value,
|
|
DenseTensor* out) {
|
|
PADDLE_ENFORCE_GT(
|
|
index_num,
|
|
0,
|
|
common::errors::InvalidArgument(
|
|
"The value 'index_num' for Op(shard_index) must be greater than 0, "
|
|
"but the value given is %d.",
|
|
index_num));
|
|
PADDLE_ENFORCE_GT(nshards,
|
|
0,
|
|
common::errors::InvalidArgument(
|
|
"The value 'nshards' for Op(shard_index) must be "
|
|
"greater than 0, but the value given is %d.",
|
|
nshards));
|
|
PADDLE_ENFORCE_GE(
|
|
shard_id,
|
|
0,
|
|
common::errors::InvalidArgument(
|
|
"The value 'shard_id' for Op(shard_index) must be greater or "
|
|
"equal to 0, but the value given is %d.",
|
|
shard_id));
|
|
PADDLE_ENFORCE_LT(
|
|
shard_id,
|
|
nshards,
|
|
common::errors::InvalidArgument(
|
|
"The value 'shard_id' for Op(shard_index) must be less than "
|
|
"nshards (%d), but the value given is %d.",
|
|
nshards,
|
|
shard_id));
|
|
|
|
out->Resize(in.dims());
|
|
out->set_lod(in.lod());
|
|
auto* in_data = in.data<T>();
|
|
auto* out_data = dev_ctx.template Alloc<T>(out);
|
|
int64_t numel = in.numel();
|
|
if (numel == 0) return;
|
|
auto stream = dev_ctx.stream();
|
|
ShardIndexInner<T>
|
|
<<<(numel + PADDLE_CUDA_NUM_THREADS - 1) / PADDLE_CUDA_NUM_THREADS,
|
|
PADDLE_CUDA_NUM_THREADS,
|
|
0,
|
|
stream>>>(
|
|
in_data, out_data, numel, index_num, nshards, shard_id, ignore_value);
|
|
}
|
|
|
|
} // namespace phi
|
|
|
|
PD_REGISTER_KERNEL(
|
|
shard_index, GPU, ALL_LAYOUT, phi::ShardIndexKernel, int, int64_t) {}
|