[MLIR] Crash in NVPTXDAGToDAGISel::getAddrSpace when lowering via NVVM pipeline

Hi, I found a crash when lowering a GPU program through the NVVM pipeline.

**Input Program**
```llvm
module {
func.func @launch_copy(%arg0: memref, %arg1: memref, %arg2: index, %arg3: index, %dyn_arg: index) -> () {
%c1 = arith.constant 1 : index
%grid_x = arith.divui %arg2, %arg3 : index
%grid_y = arith.muli %grid_x, %dyn_arg : index
%grid_z = arith.constant 1 : index
%block_x = arith.constant 512 : index
%block_y = arith.constant 2 : index
%block_z = arith.constant 1 : index
gpu.launch blocks(%bx, %by, %bz) in (%sz_bx = %grid_x, %sz_by = %grid_y, %sz_bz = %grid_z)
threads(%tx, %ty, %tz) in (%sz_tx = %block_x, %sz_ty = %block_y, %sz_tz = %block_z) {
%idx = arith.muli %by, %block_y : index
%idx_0 = arith.addi %idx, %ty : index
%c0 = arith.constant 0.0 : f32
%cond = arith.cmpf "olt", %c0, %c0 : f32
cf.cond_br %cond, ^bb1, ^bb2
^bb1:
%val = memref.load %arg0[%idx_0] : memref
memref.store %val, %arg1[%idx_0] : memref
cf.br ^bb3
^bb2:
%neg_val = arith.constant -1.0 : f32
memref.store %neg_val, %arg1[%idx_0] : memref
cf.br ^bb3
^bb3:
gpu.terminator
}
return
}
}
```

**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 '@launch_copy_kernel'
#0 0x000059545e04ed32 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) /llvm-project/llvm/lib/Support/Unix/Signals.inc:881:22
#1 0x000059545e04f1f6 PrintStackTraceSignalHandler(void*) /llvm-project/llvm/lib/Support/Unix/Signals.inc:963:1
#2 0x000059545e04c568 llvm::sys::RunSignalHandlers() /llvm-project/llvm/lib/Support/Signals.cpp:108:20
#3 0x000059545e04e5e6 SignalHandler(int, siginfo_t*, void*) /llvm-project/llvm/lib/Support/Unix/Signals.inc:448:14
#4 0x000074e1b7a45cb0 (/usr/lib/x86_64-linux-gnu/libc.so.6+0x45cb0)
#5 0x000074e1b7aa648c pthread_kill (/usr/lib/x86_64-linux-gnu/libc.so.6+0xa648c)
#6 0x000074e1b7a45b7e raise (/usr/lib/x86_64-linux-gnu/libc.so.6+0x45b7e)
#7 0x000074e1b7a288ec abort (/usr/lib/x86_64-linux-gnu/libc.so.6+0x288ec)
#8 0x000059545e026c98 bindingsErrorHandler(void*, char const*, bool) /llvm-project/llvm/lib/Support/ErrorHandling.cpp:253:55
#9 0x000059546157987a llvm::NVPTXDAGToDAGISel::getAddrSpace(llvm::MemSDNode const*) /llvm-project/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp:514:1
#10 0x000059546157cfa2 llvm::NVPTXDAGToDAGISel::tryStore(llvm::SDNode*) /llvm-project/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp:1395:42
#11 0x0000595461578132 llvm::NVPTXDAGToDAGISel::Select(llvm::SDNode*) /llvm-project/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp:115:5
#12 0x000059546506ac79 llvm::SelectionDAGISel::DoInstructionSelection() /llvm-project/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:1385:5
#13 0x0000595465069f38 llvm::SelectionDAGISel::CodeGenAndEmitDAG() /llvm-project/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:1154:3
#14 0x00005954650682fe 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
#15 0x000059546506e194 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) /llvm-project/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:1950:33
#16 0x0000595465066ab8 llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) /llvm-project/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:648:7
#17 0x0000595461577e07 llvm::NVPTXDAGToDAGISel::runOnMachineFunction(llvm::MachineFunction&) /llvm-project/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp:73:1
#18 0x0000595465065654 llvm::SelectionDAGISelLegacy::runOnMachineFunction(llvm::MachineFunction&) /llvm-project/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:395:40
#19 0x0000595465543f7c llvm::MachineFunctionPass::runOnFunction(llvm::Function&) /llvm-project/llvm/lib/CodeGen/MachineFunctionPass.cpp:94:30
#20 0x000059546c100b73 llvm::FPPassManager::runOnFunction(llvm::Function&) /llvm-project/llvm/lib/IR/LegacyPassManager.cpp:1416:20
#21 0x000059546c100eef llvm::FPPassManager::runOnModule(llvm::Module&) /llvm-project/llvm/lib/IR/LegacyPassManager.cpp:1470:13
#22 0x000059546c101468 (anonymous namespace)::MPPassManager::runOnModule(llvm::Module&) /llvm-project/llvm/lib/IR/LegacyPassManager.cpp:1561:20
#23 0x000059546c0fcb0e llvm::legacy::PassManagerImpl::run(llvm::Module&) /llvm-project/llvm/lib/IR/LegacyPassManager.cpp:533:13
#24 0x000059546c101de5 llvm::legacy::PassManager::run(llvm::Module&) /llvm-project/llvm/lib/IR/LegacyPassManager.cpp:1698:1
#25 0x0000595463e905a9 mlir::LLVM::ModuleToObject::translateModuleToISA(llvm::Module&, llvm::TargetMachine&, llvm::function_ref) /llvm-project/mlir/lib/Target/LLVM/ModuleToObject.cpp:221:3
#26 0x000059546146d5f5 (anonymous namespace)::NVPTXSerializer::moduleToObject(llvm::Module&) /llvm-project/mlir/lib/Target/LLVM/NVVM/Target.cpp:718:32
#27 0x0000595463e90b5d mlir::LLVM::ModuleToObject::run() /llvm-project/mlir/lib/Target/LLVM/ModuleToObject.cpp:281:36
#28 0x000059546146dc8e (anonymous namespace)::NVVMTargetAttrImpl::serializeToObject(mlir::Attribute, mlir::Operation*, mlir::gpu::TargetOptions const&) const /llvm-project/mlir/lib/Target/LLVM/NVVM/Target.cpp:766:8
#29 0x000059546146f72f 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
#30 0x000059546b0fb532 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
#31 0x00005954611422ac (anonymous namespace)::moduleSerializer(mlir::gpu::GPUModuleOp, mlir::gpu::OffloadingLLVMTranslationAttrInterface, mlir::gpu::TargetOptions const&) /llvm-project/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp:97:10
#32 0x00005954611426f5 mlir::gpu::transformGpuModulesToBinaries(mlir::Operation*, mlir::gpu::OffloadingLLVMTranslationAttrInterface, mlir::gpu::TargetOptions const&) /llvm-project/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp:130:19
#33 0x0000595461141fe4 (anonymous namespace)::GpuModuleToBinaryPass::runOnOperation() /llvm-project/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp:74:13
#34 0x000059546cbdbff3 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
#35 0x000059546cbe3e96 void llvm::function_ref::callback_fn(long) /llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:47:40
#36 0x000059545e070732 llvm::function_ref::operator()() const /llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:69:62
#37 0x000059546cbe9a2d void mlir::MLIRContext::executeAction(llvm::function_ref, llvm::ArrayRef, mlir::Pass&) /llvm-project/mlir/include/mlir/IR/MLIRContext.h:297:3
#38 0x000059546cbdc67b mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /llvm-project/mlir/lib/Pass/Pass.cpp:619:23
#39 0x000059546cbdcce8 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
#40 0x000059546cbe2869 mlir::PassManager::runPasses(mlir::Operation*, mlir::AnalysisManager) /llvm-project/mlir/lib/Pass/Pass.cpp:1128:40
#41 0x000059546cbe236f mlir::PassManager::run(mlir::Operation*) /llvm-project/mlir/lib/Pass/Pass.cpp:1102:69
#42 0x000059545e158e8c performActions(llvm::raw_ostream&, std::shared_ptr const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) /llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:596:13
#43 0x000059545e1596d2 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
#44 0x000059545e159ce1 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
#45 0x000059545e15aea5 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
#46 0x000059546d141d15 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
#47 0x000059546d1413de 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
#48 0x000059545e159fab mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) /llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:789:47
#49 0x000059545e15a2db mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) /llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:835:13
#50 0x000059545e15a499 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) /llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:851:21
#51 0x000059545e022901 main /llvm-project/mlir/tools/mlir-opt/mlir-opt.cpp:347:32
#52 0x000074e1b7a2a601 (/usr/lib/x86_64-linux-gnu/libc.so.6+0x2a601)
#53 0x000074e1b7a2a718 __libc_start_main (/usr/lib/x86_64-linux-gnu/libc.so.6+0x2a718)
#54 0x000059545e022435 _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

Reproduce the crash with input.mlir using mlir-opt input.mlir -gpu-lower-to-nvvm-pipeline, then inspect llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp, especially getAddrSpace and tryStore. Done means the supplied MLIR input no longer aborts during NVPTX instruction selection and a regression test covers the failure.

Written by the indexing model from the issue text.