brintos / llvm-project-archived public Read only

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