[PATCH] Removing unnecessary addrspacecasts from non-generic address spaces to the generic address space

[Justin Holewinski]

  Other than the inline comments I made, this looks good to me!

  Overall, I would like to see this become a proper IR pass in the 'opt' pipeline.  Having this optimization earlier may expose more optimization opportunities in the transformed code.


================
Comment at: lib/Target/NVPTX/NVPTXFavorNonGenericAddrSpaces.cpp:99
@@ +98,3 @@
+  // generic address space.
+  return (SrcTy->getAddressSpace() != NVPTX::PTXLdStInstCode::GENERIC &&
+          DestTy->getAddressSpace() == NVPTX::PTXLdStInstCode::GENERIC);
----------------
The correct enum to use here is actually AddressSpace::ADDRESS_SPACE_GENERIC from NVPTXBaseInfo.h.  The PTXLdStInstCode is to track the address space within instructions in the MachineInstr encoding.  They happen to match here, but its not guaranteed.

================
Comment at: test/CodeGen/NVPTX/addrspacecast.ll:7
@@ -6,4 +6,3 @@
 ; PTX32: conv1
-; PTX32: cvta.global.u32
-; PTX32: ld.u32
+; PTX32: ld.global.u32
 ; PTX64: conv1
----------------
Can we add an option to disable the addrspace conversion pass?  The purpose of this test is to ensure that addrspacecast SDNodes are properly selected.


http://llvm-reviews.chandlerc.com/D3235