1 /* Copyright 2016 The TensorFlow Authors. All Rights Reserved. 2 3 Licensed under the Apache License, Version 2.0 (the "License"); 4 you may not use this file except in compliance with the License. 5 You may obtain a copy of the License at 6 7 http://www.apache.org/licenses/LICENSE-2.0 8 9 Unless required by applicable law or agreed to in writing, software 10 distributed under the License is distributed on an "AS IS" BASIS, 11 WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. 12 See the License for the specific language governing permissions and 13 limitations under the License. 14 ==============================================================================*/ 15 16 #if GOOGLE_CUDA 17 18 #define EIGEN_USE_GPU 19 20 #include "tensorflow/core/kernels/multinomial_op.h" 21 22 #include <assert.h> 23 #include <stdio.h> 24 25 #include "third_party/eigen3/unsupported/Eigen/CXX11/Tensor" 26 #include "tensorflow/core/framework/tensor_types.h" 27 #include "tensorflow/core/kernels/random_op.h" 28 #include "tensorflow/core/lib/random/philox_random.h" 29 #include "tensorflow/core/lib/random/random_distributions.h" 30 #include "tensorflow/core/util/cuda_kernel_helper.h" 31 32 namespace tensorflow { 33 34 namespace functor { 35 36 using GPUDevice = Eigen::GpuDevice; 37 38 // Kernel for Multinomial op. Data is interpreted to have the following shapes: 39 // scores: [B, S, C]; maxima: [B, S]; output: [B, S]. 40 template <typename OutputType> 41 __global__ void MultinomialKernel(int32 nthreads, const int32 num_classes, 42 const int32 num_samples, const float* scores, 43 const float* maxima, OutputType* output) { 44 CUDA_1D_KERNEL_LOOP(index, nthreads) { 45 const int maxima_idx = index / num_classes; 46 if (ldg(maxima + maxima_idx) == ldg(scores + index)) { 47 using UnsignedOutputType = typename std::make_unsigned<OutputType>::type; 48 CudaAtomicMax(reinterpret_cast<UnsignedOutputType*>(output + maxima_idx), 49 static_cast<UnsignedOutputType>(index % num_classes)); 50 } 51 } 52 } 53 54 template <typename T, typename OutputType> 55 struct MultinomialFunctor<GPUDevice, T, OutputType> { 56 void operator()(OpKernelContext* ctx, const GPUDevice& d, 57 typename TTypes<T>::ConstMatrix logits, 58 typename TTypes<float>::Flat noises, 59 typename TTypes<float>::Flat scores, 60 typename TTypes<float>::Flat maxima, int batch_size, 61 int num_classes, int num_samples, 62 const random::PhiloxRandom& gen, 63 typename TTypes<OutputType>::Matrix output) { 64 // Uniform, [0, 1). 65 typedef random::UniformDistribution<random::PhiloxRandom, float> Dist; 66 functor::FillPhiloxRandom<GPUDevice, Dist>()(ctx, d, gen, noises.data(), 67 noises.size(), Dist()); 68 69 #if defined(EIGEN_HAS_INDEX_LIST) 70 Eigen::IndexList<Eigen::type2index<2>> kTwo; 71 Eigen::IndexList<int, int, int> bsc; 72 bsc.set(0, batch_size); 73 bsc.set(1, num_samples); 74 bsc.set(2, num_classes); 75 76 Eigen::IndexList<int, Eigen::type2index<1>, int> boc; 77 boc.set(0, batch_size); 78 boc.set(2, num_classes); 79 80 Eigen::IndexList<Eigen::type2index<1>, int, Eigen::type2index<1>> oso; 81 oso.set(1, num_samples); 82 #else 83 Eigen::array<int, 1> kTwo{2}; 84 Eigen::array<int, 3> bsc{batch_size, num_samples, num_classes}; 85 Eigen::array<int, 3> boc{batch_size, 1, num_classes}; 86 Eigen::array<int, 3> oso{1, num_samples, 1}; 87 #endif 88 89 // Calculates "scores = logits - log(-log(noises))"; B*C*S elements. 90 // NOTE: we don't store back to "noises" because having it appear on both 91 // sides is potentially unsafe (e.g. Eigen may use ldg() to load RHS data). 92 // 2e-30 is chosen so as to be small enough to only change 0 -> 2e-30 while 93 // not affect any of the other numbers (smallest is ~1e-7), but not so small 94 // that log(x) == -inf, which is why it needs to be larger than 0 in the 95 // first place. 96 To32Bit(scores).device(d) = 97 To32Bit(logits).reshape(boc).broadcast(oso).template cast<float>() - 98 ((-((To32Bit(noises) + 2e-30f).log())).log()); 99 100 // Max-reduce along classes for each (batch, sample). 101 To32Bit(maxima).device(d) = To32Bit(scores).reshape(bsc).maximum(kTwo); 102 103 // Necessary for atomicMax() inside the kernel. 104 output.device(d) = output.constant(0LL); 105 106 const int32 work_items = batch_size * num_samples * num_classes; 107 CudaLaunchConfig config = GetCudaLaunchConfig(work_items, d); 108 MultinomialKernel<<<config.block_count, config.thread_per_block, 0, 109 d.stream()>>>(config.virtual_thread_count, num_classes, 110 num_samples, scores.data(), maxima.data(), 111 output.data()); 112 } 113 }; 114 115 // Explicit instantiation of the GPU functors. 116 template struct MultinomialFunctor<GPUDevice, Eigen::half, int32>; 117 template struct MultinomialFunctor<GPUDevice, float, int32>; 118 template struct MultinomialFunctor<GPUDevice, double, int32>; 119 template struct MultinomialFunctor<GPUDevice, int32, int32>; 120 template struct MultinomialFunctor<GPUDevice, int64, int32>; 121 122 template struct MultinomialFunctor<GPUDevice, Eigen::half, int64>; 123 template struct MultinomialFunctor<GPUDevice, float, int64>; 124 template struct MultinomialFunctor<GPUDevice, double, int64>; 125 template struct MultinomialFunctor<GPUDevice, int32, int64>; 126 template struct MultinomialFunctor<GPUDevice, int64, int64>; 127 128 } // namespace functor 129 } // namespace tensorflow 130 131 #endif // GOOGLE_CUDA 132
