[llvm] [NVPTX] Lower invalid `ISD::ADDRSPACECAST` (PR #125607)

[Matt Arsenault via llvm-commits]

================
@@ -99,6 +99,20 @@ define i32 @conv8(ptr %ptr) {
   ret i32 %val
 }
 
+; ALL-LABEL: conv9
+define i32 @conv9(ptr addrspace(1) %ptr) {
+; CLS32: cvta.global.u32
+; CLS32: cvta.to.shared.u32
+; CLS64: cvta.global.u64
+; CLS64: cvta.to.shared.u64
+; PTRCONV: cvt.u32.u64
+; NOPTRCONV-NOT: cvt.u32.u64
+; ALL: ld.shared.u32
+  %specptr = addrspacecast ptr addrspace(1) %ptr to ptr addrspace(3)
----------------
arsenm wrote:

Done in #127487, this happened to come up again since 64-bit flat atomicrmws have apparently been broken since October at -O0

https://github.com/llvm/llvm-project/pull/125607