github.com/rohankumardubey/aresdb@v0.0.2-0.20190517170215-e54e3ca06b9c/query/transform.hpp (about) 1 // Copyright (c) 2017-2018 Uber Technologies, Inc. 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 #ifndef QUERY_TRANSFORM_HPP_ 16 #define QUERY_TRANSFORM_HPP_ 17 18 #include <cuda_runtime.h> 19 #include <thrust/device_vector.h> 20 #include <thrust/execution_policy.h> 21 #include <thrust/host_vector.h> 22 #include <thrust/system/cuda/execution_policy.h> 23 #include <algorithm> 24 #include <cfloat> 25 #include <cstdint> 26 #include <vector> 27 #include "query/algorithm.hpp" 28 #include "query/binder.hpp" 29 #include "query/functor.hpp" 30 #include "query/iterator.hpp" 31 #include "query/time_series_aggregate.h" 32 #include "query/utils.hpp" 33 34 void CheckCUDAError(const char *message); 35 36 namespace ares { 37 38 template<typename OutputIterator, typename FunctorType> 39 class TransformContext { 40 public: 41 TransformContext( 42 OutputIterator outputIter, 43 int indexVectorLength, 44 FunctorType functorType, 45 void *cudaStream) 46 : outputIter(outputIter), 47 indexVectorLength(indexVectorLength), 48 functorType(functorType), 49 cudaStream(reinterpret_cast<cudaStream_t>(cudaStream)) {} 50 51 public: 52 cudaStream_t getStream() const { 53 return cudaStream; 54 } 55 56 template<typename InputIterator> 57 int run(uint32_t *indexVector, InputIterator inputIter) { 58 typedef typename InputIterator::value_type::head_type InputValueType; 59 typedef typename OutputIterator::value_type::head_type OutputValueType; 60 61 UnaryFunctor<OutputValueType, InputValueType> f(functorType); 62 return thrust::transform(GET_EXECUTION_POLICY(cudaStream), inputIter, 63 inputIter + indexVectorLength, outputIter, f) - outputIter; 64 } 65 66 template<typename LHSIterator, typename RHSIterator> 67 int run(uint32_t *indexVector, LHSIterator lhsIter, RHSIterator rhsIter) { 68 typedef typename common_type< 69 typename LHSIterator::value_type::head_type, 70 typename RHSIterator::value_type::head_type>::type InputValueType; 71 72 typedef typename OutputIterator::value_type::head_type OutputValueType; 73 74 BinaryFunctor<OutputValueType, InputValueType> f(functorType); 75 76 return thrust::transform(GET_EXECUTION_POLICY(cudaStream), lhsIter, 77 lhsIter + indexVectorLength, rhsIter, outputIter, f) - outputIter; 78 } 79 80 protected: 81 OutputIterator outputIter; 82 int indexVectorLength; 83 FunctorType functorType; 84 cudaStream_t cudaStream; 85 }; 86 87 // OutputVectorBinder bind a OutputVector to a output iterator. 88 template<int NInput, typename FunctorType> 89 class OutputVectorBinder { 90 public: 91 explicit OutputVectorBinder(OutputVector outputVector, 92 std::vector<InputVector> inputVectors, 93 uint32_t *indexVector, 94 int indexVectorLength, 95 uint32_t *baseCounts, 96 uint32_t startCount, 97 FunctorType functorType, 98 void *cudaStream) : 99 output(outputVector), 100 inputs(inputVectors), 101 indexVector(indexVector), 102 indexVectorLength(indexVectorLength), 103 baseCounts(baseCounts), 104 startCount(startCount), 105 functorType(functorType), 106 cudaStream(cudaStream) {} 107 108 int bind() { 109 switch (output.Type) { 110 case ScratchSpaceOutput: 111 return transformScratchSpaceOutput(output.Vector.ScratchSpace); 112 case MeasureOutput: 113 return transformMeasureOutput(output.Vector.Measure); 114 case DimensionOutput: 115 return transformDimensionOutput(output.Vector.Dimension); 116 default:throw std::invalid_argument("Unsupported output vector type"); 117 } 118 } 119 120 // Shared by all output iterator; 121 template<typename OutputIterator> 122 int transform(OutputIterator outputIter) { 123 typedef TransformContext<OutputIterator, FunctorType> Context; 124 Context ctx( 125 outputIter, indexVectorLength, functorType, cudaStream); 126 InputVectorBinder<Context, NInput> binder( 127 ctx, inputs, indexVector, baseCounts, startCount); 128 return binder.bind(); 129 } 130 131 private: 132 OutputVector output; 133 std::vector<InputVector> inputs; 134 uint32_t *indexVector; 135 int indexVectorLength; 136 uint32_t *baseCounts; 137 uint32_t startCount; 138 FunctorType functorType; 139 void *cudaStream; 140 141 // Declaration of transform functions for different type of output iterators. 142 // Their actual definition is in different cu files so that we can compile 143 // them in parallel. 144 int transformDimensionOutput(DimensionOutputVector output); 145 int transformMeasureOutput(MeasureOutputVector output); 146 int transformScratchSpaceOutput(ScratchSpaceVector output); 147 }; 148 149 } // namespace ares 150 #endif // QUERY_TRANSFORM_HPP_