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
14using namespace mlir;
15using namespace mlir::bufferization;
16
17namespace {
18///
19struct GPUTerminatorOpInterface
20 : public BufferDeallocationOpInterface::ExternalModel<
21 GPUTerminatorOpInterface, gpu::TerminatorOp> {
22 FailureOr<Operation *> process(Operation *op, DeallocationState &state,
23 const DeallocationOptions &options) const {
24 SmallVector<Value> updatedOperandOwnerships;
25 return deallocation_impl::insertDeallocOpForReturnLike(
26 state, op, operands: {}, updatedOperandOwnerships);
27 }
28};
29
30} // namespace
31
32void mlir::gpu::registerBufferDeallocationOpInterfaceExternalModels(
33 DialectRegistry &registry) {
34 registry.addExtension(extensionFn: +[](MLIRContext *ctx, GPUDialect *dialect) {
35 gpu::TerminatorOp::attachInterface<GPUTerminatorOpInterface>(*ctx);
36 });
37}
38

source code of mlir/lib/Dialect/GPU/Transforms/BufferDeallocationOpInterfaceImpl.cpp