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 "external/cub_archive/cub/device/device_histogram.cuh" 21 #include "tensorflow/core/framework/op_kernel.h" 22 #include "tensorflow/core/framework/register_types.h" 23 #include "tensorflow/core/framework/tensor.h" 24 #include "tensorflow/core/framework/tensor_shape.h" 25 #include "tensorflow/core/kernels/bincount_op.h" 26 #include "tensorflow/core/platform/logging.h" 27 #include "tensorflow/core/platform/types.h" 28 #include "tensorflow/core/util/cuda_kernel_helper.h" 29 30 namespace tensorflow { 31 32 typedef Eigen::GpuDevice GPUDevice; 33 34 namespace functor { 35 36 template <typename T> 37 struct BincountFunctor<GPUDevice, T> { 38 static Status Compute(OpKernelContext* context, 39 const typename TTypes<int32, 1>::ConstTensor& arr, 40 const typename TTypes<T, 1>::ConstTensor& weights, 41 typename TTypes<T, 1>::Tensor& output) { 42 if (weights.size() != 0) { 43 return errors::InvalidArgument( 44 "Weights should not be passed as it should be " 45 "handled by unsorted_segment_sum"); 46 } 47 if (output.size() == 0) { 48 return Status::OK(); 49 } 50 // In case weight.size() == 0, use CUB 51 size_t temp_storage_bytes = 0; 52 const int32* d_samples = arr.data(); 53 T* d_histogram = output.data(); 54 int num_levels = output.size() + 1; 55 int32 lower_level = 0; 56 int32 upper_level = output.size(); 57 int num_samples = arr.size(); 58 const cudaStream_t& stream = GetCudaStream(context); 59 60 // The first HistogramEven is to obtain the temp storage size required 61 // with d_temp_storage = NULL passed to the call. 62 auto err = cub::DeviceHistogram::HistogramEven( 63 /* d_temp_storage */ NULL, 64 /* temp_storage_bytes */ temp_storage_bytes, 65 /* d_samples */ d_samples, 66 /* d_histogram */ d_histogram, 67 /* num_levels */ num_levels, 68 /* lower_level */ lower_level, 69 /* upper_level */ upper_level, 70 /* num_samples */ num_samples, 71 /* stream */ stream); 72 if (err != cudaSuccess) { 73 return errors::Internal( 74 "Could not launch HistogramEven to get temp storage: ", 75 cudaGetErrorString(err), "."); 76 } 77 Tensor temp_storage; 78 TF_RETURN_IF_ERROR(context->allocate_temp( 79 DataTypeToEnum<int8>::value, 80 TensorShape({static_cast<int64>(temp_storage_bytes)}), &temp_storage)); 81 82 void* d_temp_storage = temp_storage.flat<int8>().data(); 83 // The second HistogramEven is to actual run with d_temp_storage 84 // allocated with temp_storage_bytes. 85 err = cub::DeviceHistogram::HistogramEven( 86 /* d_temp_storage */ d_temp_storage, 87 /* temp_storage_bytes */ temp_storage_bytes, 88 /* d_samples */ d_samples, 89 /* d_histogram */ d_histogram, 90 /* num_levels */ num_levels, 91 /* lower_level */ lower_level, 92 /* upper_level */ upper_level, 93 /* num_samples */ num_samples, 94 /* stream */ stream); 95 if (err != cudaSuccess) { 96 return errors::Internal( 97 "Could not launch HistogramEven: ", cudaGetErrorString(err), "."); 98 } 99 return Status::OK(); 100 } 101 }; 102 103 } // end namespace functor 104 105 #define REGISTER_GPU_SPEC(type) \ 106 template struct functor::BincountFunctor<GPUDevice, type>; 107 108 TF_CALL_int32(REGISTER_GPU_SPEC); 109 TF_CALL_float(REGISTER_GPU_SPEC); 110 #undef REGISTER_GPU_SPEC 111 112 } // namespace tensorflow 113 114 #endif // GOOGLE_CUDA 115
