/* Copyright (c) 2020 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 #include #include #include #include "paddle/fluid/framework/eigen.h" #include "paddle/fluid/framework/op_registry.h" #include "paddle/fluid/framework/operator.h" #include "paddle/fluid/operators/multinomial_op.h" #include "paddle/fluid/platform/transform.h" namespace paddle { namespace operators { /* template using EigenVector = framework::EigenVector; template using EigenMatrix = framework::EigenMatrix; */ /* template __global__ void SumArrayCUDAKernel(T **in, T *out, size_t in_size) { int id = blockIdx.x * blockDim.x + threadIdx.x; // T total(read_dst ? out[id] : static_cast(0)); T total(static_cast(0)) for (int i = 0; i < in_size; ++i) { const T *tmp = in[i]; if (tmp) { total += tmp[id]; } } out[id] = total; id += blockDim.x * gridDim.x; }*/ /* template __global__ void NormalizeProbability(T* probs, int64_t rows, int64_t cols) { extern __shared__ std::vector sum_rows(rows); T val; for (int64_t i = blockId.x; i < rows; i += gridDim.x) { T sum = static_cast(0); for (int64_t j = threadIdx.x; j < cols; j += blockDim.x) { val = probs[i * cols + j]; sum += val; } } }*/ template __global__ void NormalizeProbability(T* norm_probs, const T* in_data, T* sum_rows) { // int id = blockIdx.x * blockDim.x + threadIdx.x; int id = threadIdx.x; norm_probs[id] = in_data[id] / sum_rows[0]; } template struct RandomGeneratorCudaFunctor { unsigned int seed_; __host__ __device__ RandomGeneratorCudaFunctor(int seed) : seed_(seed) {} __host__ __device__ T operator()(const unsigned int n) const { thrust::minstd_rand rng; rng.seed(seed_); thrust::uniform_real_distribution dist(0.0, 1.0); rng.discard(n); return dist(rng); } }; /* template class MultinomialCudaFunctor(T* out_data, const T* in_data, const int64_t num_samples, const bool replacement, const int64_t num_categories, const int64_t num_distributions) { }*/ template __device__ int binarySearchForMultinomial(T* cumdist, T* dist, int size, T val) { int start = 0; int end = size; // cumdist[size - 1] = 0 => all zero prob dist // CUDA_KERNEL_ASSERT(cumdist[size - 1] > static_cast(0)); while (end - start > 0) { int mid = start + (end - start) / 2; T midVal = cumdist[mid]; if (midVal < val) { start = mid + 1; } else { end = mid; } } if (start == size) { // No probability mass or precision problems; just return the // first non-zero element by setting start to size-1 here, // the code below will move it to the last non-zero probability // this actually can happen when the random number is 1 // (github pytorch issue #4858). start = size - 1; } while (start >= 1 && dist[start] == 0) start--; return start; } template __global__ void sampleMultinomialWithReplacement( T* rng, const int64_t totalSamples, T* dest, const int64_t distributions, const int64_t categories, T* normDistPrefixSum, T* normDist) { // At the moment, each warp computes one sample value in the binary // search due to divergence. It seems possible to compute multiple // values and limit divergence though later on. // global index formula for 2D grid of 1D blocks // int idx = blockIdx.y * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + // threadIdx.x; int idx = blockIdx.x * blockDim.x + threadIdx.x; for (int sample = blockIdx.x * blockDim.x + threadIdx.x; sample < totalSamples; sample += blockDim.x * gridDim.x) { // we are losing 3 out of 4 generated numbers but it's ok // this kernel is not very efficient anyway // T uniform_random = dist(rng); T uniform_random = rng[sample]; // Find the bucket that a uniform sample lies in int choice = binarySearchForMultinomial(normDistPrefixSum, normDist, categories, uniform_random); dest[sample] = choice; } } template class MultinomialOpKernel : public framework::OpKernel { public: void Compute(const framework::ExecutionContext& ctx) const override { const auto x = ctx.Input("X"); auto out = ctx.Output("Out"); const int64_t num_samples = ctx.Attr("num_samples"); const bool replacement = ctx.Attr("replacement"); auto* in_data = x->data(); auto* out_data = out->mutable_data(ctx.GetPlace()); auto in_dims = x->dims(); int64_t in_rank = in_dims.size(); const int64_t num_categories = in_dims[in_rank - 1]; const int64_t num_distributions = in_rank > 1 ? in_dims[in_rank - 2] : 1; // std::vector sum_rows(num_distributions); // SumArrayCUDAKernel(in_data, sum_rows,) VLOG(3) << "Print num_distributions " << num_distributions << "\n"; VLOG(3) << "Print num_categories " << num_categories << "\n"; VLOG(3) << "Print in_rank " << in_rank << "\n"; framework::Tensor sum_rows_t; auto* sum_rows_data = sum_rows_t.mutable_data({1}, ctx.GetPlace()); // auto* sum_rows_data = // sum_rows_t->mutable_data(framework::make_ddim({1}), ctx.GetPlace()); auto& place = *ctx.template device_context() .eigen_device(); auto eigen_input = framework::EigenVector::Flatten(*x); // auto eigen_sum_rows = framework::EigenVector::From(sum_rows_t); auto eigen_sum_rows = framework::EigenScalar::From(sum_rows_t); eigen_sum_rows.device(place) = eigen_input.sum(Eigen::DSizes(0)) .eval() .reshape(Eigen::DSizes(sum_rows_t.dims()[0])); // eigen_sum_rows.device(place) = // eigen_input.sum().eval().reshape(Eigen::DSizes(1)); dim3 grid(num_distributions); dim3 block(num_categories); // std::vector in_data_norm(num_categories); framework::Tensor norm_probs_t; auto* norm_probs_data = norm_probs_t.mutable_data({num_categories}, ctx.GetPlace()); NormalizeProbability< T><<>>( norm_probs_data, in_data, sum_rows_data); // num_distributions can only be 1. // std::vector cumulative_probs(num_categories); framework::Tensor cumulative_probs_t; auto* cumulative_probs = cumulative_probs_t.mutable_data({num_categories}, ctx.GetPlace()); // T cumulative_probs[num_categories]; int64_t size = num_categories; thrust::inclusive_scan(thrust::device, norm_probs_data, norm_probs_data + num_categories, cumulative_probs); if (replacement) { dim3 block(128); // int grid_y = 1; dim3 grid((num_samples - 1) / block.x + 1); /* // std::vector rng(num_samples); T rng[num_samples]; std::uniform_real_distribution dist(0, 1); auto gen_ptr = framework::DefaultCPUGenerator(); auto engine = gen_ptr->GetCPUEngine(); for (int s = 0; s < num_samples; s++) { rng[s] = dist(*engine); } */ std::random_device rd; auto seed = rd(); framework::Tensor rng_data_t; auto* rng_data = rng_data_t.mutable_data({num_samples}, ctx.GetPlace()); thrust::counting_iterator index_sequence_begin(0); platform::Transform trans; auto* context = static_cast( &ctx.device_context()); trans(*context, index_sequence_begin, index_sequence_begin + num_samples, rng_data, RandomGeneratorCudaFunctor(seed)); VLOG(3) << "Print enter\n"; // VLOG(3) << "Print size in_data " << // sizeof(in_data)/sizeof(in_data[num_categories-1]) << "\n"; // VLOG(3) << "Print norm_probs_data0 " << // sizeof(norm_probs_data[num_categories-1]) << "\n"; sampleMultinomialWithReplacement< T><<>>( rng_data, num_samples, out_data, num_distributions, num_categories, cumulative_probs, norm_probs_data); } // MultinomialCudaFunctor(out_data, in_data, num_samples, replacement, // num_categories, num_distributions); } }; } // namespace operators } // namespace paddle namespace ops = paddle::operators; namespace plat = paddle::platform; REGISTER_OP_CUDA_KERNEL( multinomial, ops::MultinomialOpKernel, ops::MultinomialOpKernel);