[MLIR] Operation reach NVPTX backend and cause instruction selection crash

Hi, I found that `mlir-opt` crashes instead of emitting a diagnostic when a GPU module containing ROCDL operations is lowered through the NVVM pipeline. Please see the following minimized reproducer.

**Code Example**
```llvm
module attributes {gpu.container_module} {
gpu.module @kernels {
gpu.func @capture() kernel {
%tx = rocdl.workitem.id.x : i32
%c0 = arith.constant 0 : i32
%cond = arith.cmpi eq, %tx, %c0 : i32
scf.if %cond {
gpu.printf "GPU: thread 0\n"
}
rocdl.barrier
gpu.printf "GPU: all threads here\n"
gpu.return
}
}
func.func @main() -> i32 {
%c1 = arith.constant 1 : index
%c64 = arith.constant 64 : index
gpu.launch_func @kernels::@capture blocks in (%c1, %c1, %c1) threads in (%c64, %c1, %c1)
%c0 = arith.constant 0 : i32
return %c0 : i32
}
}
```

**Stack Trace**
```bash
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace and instructions to reproduce the bug.
Stack dump:
0. Program arguments: mlir-opt input.mlir -gpu-lower-to-nvvm-pipeline
1. Running pass 'Function Pass Manager' on module 'LLVMDialectModule'.
2. Running pass 'NVPTX DAG->DAG Pattern Instruction Selection' on function '@capture'
#0 0x000059144b80ed32 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) /llvm-project/llvm/lib/Support/Unix/Signals.inc:881:22
#1 0x000059144b80f1f6 PrintStackTraceSignalHandler(void*) /llvm-project/llvm/lib/Support/Unix/Signals.inc:963:1
#2 0x000059144b80c568 llvm::sys::RunSignalHandlers() /llvm-project/llvm/lib/Support/Signals.cpp:108:20
#3 0x000059144b80e5e6 SignalHandler(int, siginfo_t*, void*) /llvm-project/llvm/lib/Support/Unix/Signals.inc:448:14
#4 0x00007a7f82645cb0 (/usr/lib/x86_64-linux-gnu/libc.so.6+0x45cb0)
#5 0x00007a7f826a648c pthread_kill (/usr/lib/x86_64-linux-gnu/libc.so.6+0xa648c)
#6 0x00007a7f82645b7e raise (/usr/lib/x86_64-linux-gnu/libc.so.6+0x45b7e)
#7 0x00007a7f826288ec abort (/usr/lib/x86_64-linux-gnu/libc.so.6+0x288ec)
#8 0x000059144b7e688d llvm::report_fatal_error(llvm::Twine const&, bool) /llvm-project/llvm/lib/Support/ErrorHandling.cpp:137:9
#9 0x000059145283ceb8 llvm::iterator_range, false, false>>&>())), llvm::filterDbgVars(llvm::iterator_range, false, false>>)::'lambda'(llvm::DbgRecord&), std::conditional, false, false>>&>()))>::iterator_category>, std::bidirectional_iterator_tag, std::forward_iterator_tag>::type>> llvm::make_filter_range, false, false>>&, llvm::filterDbgVars(llvm::iterator_range, false, false>>)::'lambda'(llvm::DbgRecord&)>(llvm::iterator_range, false, false>>&, llvm::filterDbgVars(llvm::iterator_range, false, false>>)::'lambda'(llvm::DbgRecord&)) /llvm-project/llvm/include/llvm/ADT/STLExtras.h:551:1
#10 0x000059145283c6b7 llvm::SelectionDAGISel::SelectCodeCommon(llvm::SDNode*, unsigned char const*, unsigned int, unsigned char const*) /llvm-project/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:4540:9
#11 0x000059144ed4319d llvm::NVPTXDAGToDAGISel::SelectCode(llvm::SDNode*) /llvm-project/build_de045d/lib/Target/NVPTX/NVPTXGenDAGISel.inc:77137:1
#12 0x000059144ed383ec llvm::NVPTXDAGToDAGISel::Select(llvm::SDNode*) /llvm-project/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp:200:1
#13 0x000059145282ac79 llvm::SelectionDAGISel::DoInstructionSelection() /llvm-project/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:1385:5
#14 0x0000591452829f38 llvm::SelectionDAGISel::CodeGenAndEmitDAG() /llvm-project/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:1154:3
#15 0x00005914528282fe llvm::SelectionDAGISel::SelectBasicBlock(llvm::ilist_iterator_w_bits, false, true>, llvm::ilist_iterator_w_bits, false, true>, bool&) /llvm-project/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:909:1
#16 0x000059145282e194 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) /llvm-project/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:1950:33
#17 0x0000591452826ab8 llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) /llvm-project/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:648:7
#18 0x000059144ed37e07 llvm::NVPTXDAGToDAGISel::runOnMachineFunction(llvm::MachineFunction&) /llvm-project/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp:73:1
#19 0x0000591452825654 llvm::SelectionDAGISelLegacy::runOnMachineFunction(llvm::MachineFunction&) /llvm-project/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:395:40
#20 0x0000591452d03f7c llvm::MachineFunctionPass::runOnFunction(llvm::Function&) /llvm-project/llvm/lib/CodeGen/MachineFunctionPass.cpp:94:30
#21 0x00005914598c0b73 llvm::FPPassManager::runOnFunction(llvm::Function&) /llvm-project/llvm/lib/IR/LegacyPassManager.cpp:1416:20
#22 0x00005914598c0eef llvm::FPPassManager::runOnModule(llvm::Module&) /llvm-project/llvm/lib/IR/LegacyPassManager.cpp:1470:13
#23 0x00005914598c1468 (anonymous namespace)::MPPassManager::runOnModule(llvm::Module&) /llvm-project/llvm/lib/IR/LegacyPassManager.cpp:1561:20
#24 0x00005914598bcb0e llvm::legacy::PassManagerImpl::run(llvm::Module&) /llvm-project/llvm/lib/IR/LegacyPassManager.cpp:533:13
#25 0x00005914598c1de5 llvm::legacy::PassManager::run(llvm::Module&) /llvm-project/llvm/lib/IR/LegacyPassManager.cpp:1698:1
#26 0x00005914516505a9 mlir::LLVM::ModuleToObject::translateModuleToISA(llvm::Module&, llvm::TargetMachine&, llvm::function_ref) /llvm-project/mlir/lib/Target/LLVM/ModuleToObject.cpp:221:3
#27 0x000059144ec2d5f5 (anonymous namespace)::NVPTXSerializer::moduleToObject(llvm::Module&) /llvm-project/mlir/lib/Target/LLVM/NVVM/Target.cpp:718:32
#28 0x0000591451650b5d mlir::LLVM::ModuleToObject::run() /llvm-project/mlir/lib/Target/LLVM/ModuleToObject.cpp:281:36
#29 0x000059144ec2dc8e (anonymous namespace)::NVVMTargetAttrImpl::serializeToObject(mlir::Attribute, mlir::Operation*, mlir::gpu::TargetOptions const&) const /llvm-project/mlir/lib/Target/LLVM/NVVM/Target.cpp:766:8
#30 0x000059144ec2f72f mlir::gpu::detail::TargetAttrInterfaceInterfaceTraits::FallbackModel<(anonymous namespace)::NVVMTargetAttrImpl>::serializeToObject(mlir::gpu::detail::TargetAttrInterfaceInterfaceTraits::Concept const*, mlir::Attribute, mlir::Operation*, mlir::gpu::TargetOptions const&) /llvm-project/build_de045d/tools/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.h.inc:250:1
#31 0x00005914588bb532 mlir::gpu::TargetAttrInterface::serializeToObject(mlir::Operation*, mlir::gpu::TargetOptions const&) const /llvm-project/build_de045d/tools/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.cpp.inc:20:3
#32 0x000059144e9022ac (anonymous namespace)::moduleSerializer(mlir::gpu::GPUModuleOp, mlir::gpu::OffloadingLLVMTranslationAttrInterface, mlir::gpu::TargetOptions const&) /llvm-project/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp:97:10
#33 0x000059144e9026f5 mlir::gpu::transformGpuModulesToBinaries(mlir::Operation*, mlir::gpu::OffloadingLLVMTranslationAttrInterface, mlir::gpu::TargetOptions const&) /llvm-project/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp:130:19
#34 0x000059144e901fe4 (anonymous namespace)::GpuModuleToBinaryPass::runOnOperation() /llvm-project/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp:74:13
#35 0x000059145a39bff3 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::'lambda'()::operator()() const /llvm-project/mlir/lib/Pass/Pass.cpp:613:24
#36 0x000059145a3a3e96 void llvm::function_ref::callback_fn(long) /llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:47:40
#37 0x000059144b830732 llvm::function_ref::operator()() const /llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:69:62
#38 0x000059145a3a9a2d void mlir::MLIRContext::executeAction(llvm::function_ref, llvm::ArrayRef, mlir::Pass&) /llvm-project/mlir/include/mlir/IR/MLIRContext.h:297:3
#39 0x000059145a39c67b mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /llvm-project/mlir/lib/Pass/Pass.cpp:619:23
#40 0x000059145a39cce8 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) /llvm-project/mlir/lib/Pass/Pass.cpp:688:15
#41 0x000059145a3a2869 mlir::PassManager::runPasses(mlir::Operation*, mlir::AnalysisManager) /llvm-project/mlir/lib/Pass/Pass.cpp:1128:40
#42 0x000059145a3a236f mlir::PassManager::run(mlir::Operation*) /llvm-project/mlir/lib/Pass/Pass.cpp:1102:69
#43 0x000059144b918e8c performActions(llvm::raw_ostream&, std::shared_ptr const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) /llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:596:13
#44 0x000059144b9196d2 processBuffer(llvm::raw_ostream&, std::unique_ptr>, llvm::MemoryBufferRef, mlir::MlirOptMainConfig const&, mlir::DialectRegistry&, mlir::SourceMgrDiagnosticVerifierHandler*, llvm::ThreadPoolInterface*) /llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:686:26
#45 0x000059144b919ce1 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&)::'lambda'(std::unique_ptr>, llvm::MemoryBufferRef, llvm::raw_ostream&)::operator()(std::unique_ptr>, llvm::MemoryBufferRef, llvm::raw_ostream&) const /llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:784:25
#46 0x000059144b91aea5 llvm::LogicalResult llvm::function_ref>, llvm::MemoryBufferRef const&, llvm::raw_ostream&)>::callback_fn>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&)::'lambda'(std::unique_ptr>, llvm::MemoryBufferRef, llvm::raw_ostream&)>(long, std::unique_ptr>, llvm::MemoryBufferRef const&, llvm::raw_ostream&) /llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:46:52
#47 0x000059145a901d15 llvm::function_ref>, llvm::MemoryBufferRef const&, llvm::raw_ostream&)>::operator()(std::unique_ptr>, llvm::MemoryBufferRef const&, llvm::raw_ostream&) const /llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:69:12
#48 0x000059145a9013de mlir::splitAndProcessBuffer(std::unique_ptr>, llvm::function_ref>, llvm::MemoryBufferRef const&, llvm::raw_ostream&)>, llvm::raw_ostream&, llvm::StringRef, llvm::StringRef) /llvm-project/mlir/lib/Support/ToolUtilities.cpp:30:30
#49 0x000059144b919fab mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) /llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:789:47
#50 0x000059144b91a2db mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) /llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:835:13
#51 0x000059144b91a499 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) /llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:851:21
#52 0x000059144b7e2901 main /llvm-project/mlir/tools/mlir-opt/mlir-opt.cpp:347:32
#53 0x00007a7f8262a601 (/usr/lib/x86_64-linux-gnu/libc.so.6+0x2a601)
#54 0x00007a7f8262a718 __libc_start_main (/usr/lib/x86_64-linux-gnu/libc.so.6+0x2a718)
#55 0x000059144b7e2435 _start (/mnt/raid/llvm-project/build_de045d/bin/mlir-opt+0x464c435)
Aborted mlir-opt input.mlir -gpu-lower-to-nvvm-pipeline
```

**Used Command**
```bash
mlir-opt input.mlir -gpu-lower-to-nvvm-pipeline
```

Version: de045d579791d978cce109f9348d20560e112d74