1 /* Copyright 2017 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 "third_party/eigen3/unsupported/Eigen/CXX11/Tensor" 21 #include "external/cub_archive/cub/device/device_histogram.cuh" 22 #include "tensorflow/core/framework/op_kernel.h" 23 #include "tensorflow/core/framework/register_types.h" 24 #include "tensorflow/core/framework/tensor.h" 25 #include "tensorflow/core/framework/tensor_shape.h" 26 #include "tensorflow/core/kernels/histogram_op.h" 27 #include "tensorflow/core/platform/logging.h" 28 #include "tensorflow/core/platform/types.h" 29 #include "tensorflow/core/util/cuda_kernel_helper.h" 30 31 namespace tensorflow { 32 33 typedef Eigen::GpuDevice GPUDevice; 34 35 namespace functor { 36 37 // TODO(yongtang) int64 of atomicAdd is not supported yet. 38 template <typename T, typename Tout> 39 struct HistogramFixedWidthFunctor<GPUDevice, T, Tout> { 40 static Status Compute(OpKernelContext* context, 41 const typename TTypes<T, 1>::ConstTensor& values, 42 const typename TTypes<T, 1>::ConstTensor& value_range, 43 int32 nbins, typename TTypes<Tout, 1>::Tensor& out) { 44 tensorflow::AllocatorAttributes pinned_allocator; 45 pinned_allocator.set_on_host(true); 46 pinned_allocator.set_gpu_compatible(true); 47 48 Tensor levels_tensor; 49 TF_RETURN_IF_ERROR(context->allocate_temp( 50 DataTypeToEnum<T>::value, TensorShape({nbins + 1}), &levels_tensor, 51 pinned_allocator)); 52 auto levels = levels_tensor.flat<T>(); 53 54 const double step = static_cast<double>(value_range(1) - value_range(0)) / 55 static_cast<double>(nbins); 56 levels(0) = std::numeric_limits<T>::lowest(); 57 for (int i = 1; i < nbins; i++) { 58 levels(i) = 59 static_cast<T>(static_cast<double>(value_range(0)) + step * i); 60 } 61 levels(nbins) = std::numeric_limits<T>::max(); 62 63 size_t temp_storage_bytes = 0; 64 const T* d_samples = values.data(); 65 Tout* d_histogram = out.data(); 66 int num_levels = levels.size(); 67 T* d_levels = levels.data(); 68 int num_samples = values.size(); 69 const cudaStream_t& stream = GetCudaStream(context); 70 71 // The first HistogramRange is to obtain the temp storage size required 72 // with d_temp_storage = NULL passed to the call. 73 auto err = cub::DeviceHistogram::HistogramRange( 74 /* d_temp_storage */ NULL, 75 /* temp_storage_bytes */ temp_storage_bytes, 76 /* d_samples */ d_samples, 77 /* d_histogram */ d_histogram, 78 /* num_levels */ num_levels, 79 /* d_levels */ d_levels, 80 /* num_samples */ num_samples, 81 /* stream */ stream); 82 if (err != cudaSuccess) { 83 return errors::Internal( 84 "Could not launch HistogramRange to get temp storage: ", 85 cudaGetErrorString(err), "."); 86 } 87 88 Tensor temp_storage; 89 TF_RETURN_IF_ERROR(context->allocate_temp( 90 DataTypeToEnum<int8>::value, 91 TensorShape({static_cast<int64>(temp_storage_bytes)}), &temp_storage)); 92 93 void* d_temp_storage = temp_storage.flat<int8>().data(); 94 95 // The second HistogramRange is to actual run with d_temp_storage 96 // allocated with temp_storage_bytes. 97 err = cub::DeviceHistogram::HistogramRange( 98 /* d_temp_storage */ d_temp_storage, 99 /* temp_storage_bytes */ temp_storage_bytes, 100 /* d_samples */ d_samples, 101 /* d_histogram */ d_histogram, 102 /* num_levels */ num_levels, 103 /* d_levels */ d_levels, 104 /* num_samples */ num_samples, 105 /* stream */ stream); 106 if (err != cudaSuccess) { 107 return errors::Internal( 108 "Could not launch HistogramRange: ", cudaGetErrorString(err), "."); 109 } 110 111 return Status::OK(); 112 } 113 }; 114 115 } // end namespace functor 116 117 #define REGISTER_GPU_SPEC(type) \ 118 template struct functor::HistogramFixedWidthFunctor<GPUDevice, type, int32>; 119 120 TF_CALL_GPU_NUMBER_TYPES(REGISTER_GPU_SPEC); 121 #undef REGISTER_GPU_SPEC 122 123 } // namespace tensorflow 124 125 #endif // GOOGLE_CUDA 126
