1 /* Copyright 2021 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 #ifndef TENSORFLOW_COMPILER_XLA_SERVICE_GPU_CUDNN_VECTORIZE_CONVOLUTIONS_H_ 17 #define TENSORFLOW_COMPILER_XLA_SERVICE_GPU_CUDNN_VECTORIZE_CONVOLUTIONS_H_ 18 19 #include <utility> 20 21 #include "tensorflow/compiler/xla/service/hlo_pass_interface.h" 22 #include "tensorflow/compiler/xla/statusor.h" 23 24 namespace xla { 25 namespace gpu { 26 27 // Changes the shape of cudnn convolutions to allow faster "vectorized" 28 // algorithms. 29 // 30 // On sm61+ will convert int8_t convolutions from 31 // 32 // - [N, C, H, W] to [N, C/4, H, W, 4], 33 // 34 // assuming C is divisible by 4. 35 // 36 // On sm75+ will convert int8_t convolutions from 37 // 38 // - [N, C, H, W] to [N, C/32, H, W, 32], 39 // - [N, C/4, H, W, 4] to [N, C/32, H, W, 32], and 40 // - [N, C, H, W] to [N, C/4, H, W, 4] (same as sm61+), 41 // 42 // assuming C is divisible by 4 or 32. 43 // 44 // This pass will not pad the channel dim to a multiple of 4 or 32, so you 45 // should run CudnnPadForConvolutions before this. 46 class CudnnVectorizeConvolutions : public HloModulePass { 47 public: CudnnVectorizeConvolutions(se::CudaComputeCapability compute_capability)48 explicit CudnnVectorizeConvolutions( 49 se::CudaComputeCapability compute_capability) 50 : compute_capability_(compute_capability) {} 51 name()52 absl::string_view name() const override { 53 return "cudnn_vectorize_convolutions"; 54 } 55 using HloPassInterface::Run; 56 StatusOr<bool> Run( 57 HloModule* module, 58 const absl::flat_hash_set<absl::string_view>& execution_threads) override; 59 60 private: 61 se::CudaComputeCapability compute_capability_; 62 }; 63 64 } // namespace gpu 65 } // namespace xla 66 67 #endif // TENSORFLOW_COMPILER_XLA_SERVICE_GPU_CUDNN_VECTORIZE_CONVOLUTIONS_H_ 68