1 //===- NVGPUToNVVMPass.h - Convert NVGPU to NVVM dialect --------*- C++ -*-===// 2 // 3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4 // See https://llvm.org/LICENSE.txt for license information. 5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6 // 7 //===----------------------------------------------------------------------===// 8 #ifndef MLIR_CONVERSION_NVGPUTONVVM_NVGPUTONVVMPASS_H_ 9 #define MLIR_CONVERSION_NVGPUTONVVM_NVGPUTONVVMPASS_H_ 10 11 #include <memory> 12 13 namespace mlir { 14 15 class Attribute; 16 class LLVMTypeConverter; 17 class MemRefType; 18 class MLIRContext; 19 class RewritePatternSet; 20 class Pass; 21 22 #define GEN_PASS_DECL_CONVERTNVGPUTONVVMPASS 23 #include "mlir/Conversion/Passes.h.inc" 24 25 namespace nvgpu { 26 class MBarrierGroupType; 27 28 /// Returns the memory space attribute of the mbarrier object. 29 Attribute getMbarrierMemorySpace(MLIRContext *context, 30 MBarrierGroupType barrierType); 31 32 /// Return the memref type that can be used to represent an mbarrier object. 33 MemRefType getMBarrierMemrefType(MLIRContext *context, 34 MBarrierGroupType barrierType); 35 } // namespace nvgpu 36 37 void populateNVGPUToNVVMConversionPatterns(const LLVMTypeConverter &converter, 38 RewritePatternSet &patterns); 39 } // namespace mlir 40 41 #endif // MLIR_CONVERSION_NVGPUTONVVM_NVGPUTONVVMPASS_H_ 42