1 //===- BufferDeallocationOpInterfaceImpl.cpp ------------------------------===// 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 9 #include "mlir/Dialect/GPU/Transforms/BufferDeallocationOpInterfaceImpl.h" 10 #include "mlir/Dialect/Bufferization/IR/BufferDeallocationOpInterface.h" 11 #include "mlir/Dialect/Bufferization/IR/Bufferization.h" 12 #include "mlir/Dialect/GPU/IR/GPUDialect.h" 13 14 using namespace mlir; 15 using namespace mlir::bufferization; 16 17 namespace { 18 /// 19 struct GPUTerminatorOpInterface 20 : public BufferDeallocationOpInterface::ExternalModel< 21 GPUTerminatorOpInterface, gpu::TerminatorOp> { process__anon28ccad050111::GPUTerminatorOpInterface22 FailureOr<Operation *> process(Operation *op, DeallocationState &state, 23 const DeallocationOptions &options) const { 24 SmallVector<Value> updatedOperandOwnerships; 25 return deallocation_impl::insertDeallocOpForReturnLike( 26 state, op, {}, updatedOperandOwnerships); 27 } 28 }; 29 30 } // namespace 31 registerBufferDeallocationOpInterfaceExternalModels(DialectRegistry & registry)32void mlir::gpu::registerBufferDeallocationOpInterfaceExternalModels( 33 DialectRegistry ®istry) { 34 registry.addExtension(+[](MLIRContext *ctx, GPUDialect *dialect) { 35 gpu::TerminatorOp::attachInterface<GPUTerminatorOpInterface>(*ctx); 36 }); 37 } 38