Use MLIR target configuration for lowering to ROCDL.

This is now exposed as the method configureGpuToROCDLConversionLegality(), so
we can easily use it in kernel_lowering.

PiperOrigin-RevId: 347814204
Change-Id: Id356177c1c6c127c331c858bbce63ef25310b93d
This commit is contained in:
Adrian Kuegel 2020-12-16 06:57:56 -08:00 committed by TensorFlower Gardener
parent 2492f10eca
commit decac0e638

View File

@ -227,17 +227,7 @@ class LowerToROCDLPass
::mlir::populateAffineToStdConversionPatterns(patterns, m.getContext());
::mlir::ConversionTarget target(getContext());
target.addIllegalDialect<::mlir::gpu::GPUDialect>();
target.addIllegalOp<mlir::LLVM::CosOp, mlir::LLVM::ExpOp,
mlir::LLVM::FAbsOp, mlir::LLVM::FCeilOp,
mlir::LLVM::LogOp, mlir::LLVM::Log10Op,
mlir::LLVM::Log2Op, mlir::LLVM::SinOp>();
target.addIllegalOp<mlir::FuncOp>();
target.addLegalDialect<::mlir::LLVM::LLVMDialect>();
target.addLegalDialect<::mlir::ROCDL::ROCDLDialect>();
// TODO(csigg): Remove once we support replacing non-root ops.
target.addLegalOp<::mlir::gpu::GPUModuleOp, ::mlir::gpu::ModuleEndOp,
::mlir::gpu::YieldOp>();
::mlir::configureGpuToROCDLConversionLegality(target);
if (failed(mlir::applyFullConversion(m, target, std::move(patterns)))) {
signalPassFailure();
}