[PATCH] [NVPTX] run SROA after NVPTXFavorNonGenericAddrSpaces

[Jingyue Wu]

Hi eliben, jholewinski,

This cleans up most allocas NVPTXLowerKernelArgs emits for byval
parameters.

http://reviews.llvm.org/D10322

Files:
  lib/Target/NVPTX/NVPTXTargetMachine.cpp
  test/CodeGen/NVPTX/bug21465.ll

Index: lib/Target/NVPTX/NVPTXTargetMachine.cpp
===================================================================
--- lib/Target/NVPTX/NVPTXTargetMachine.cpp
+++ lib/Target/NVPTX/NVPTXTargetMachine.cpp
@@ -165,6 +165,11 @@
   addPass(createGenericToNVVMPass());
   addPass(createNVPTXLowerKernelArgsPass(&getNVPTXTargetMachine()));
   addPass(createNVPTXFavorNonGenericAddrSpacesPass());
+  // NVPTXLowerKernelArgs emits alloca for byval parameters which can often
+  // be eliminated by SROA. We do not run SROA right after NVPTXLowerKernelArgs
+  // because we plan to merge NVPTXLowerKernelArgs and
+  // NVPTXFavorNonGenericAddrSpaces into one pass.
+  addPass(createSROAPass());
   // FavorNonGenericAddrSpaces shortcuts unnecessary addrspacecasts, and leave
   // them unused. We could remove dead code in an ad-hoc manner, but that
   // requires manual work and might be error-prone.
Index: test/CodeGen/NVPTX/bug21465.ll
===================================================================
--- test/CodeGen/NVPTX/bug21465.ll
+++ test/CodeGen/NVPTX/bug21465.ll
@@ -14,9 +14,10 @@
 ; CHECK: addrspacecast %struct.S* %input to %struct.S addrspace(101)*
   %b = getelementptr inbounds %struct.S, %struct.S* %input, i64 0, i32 1
   %0 = load i32, i32* %b, align 4
-; PTX: ld.param.u32 %r{{[0-9]+}}, {{\[}}[[BASE:%rd[0-9]+]]{{\]}}
-; PTX-NEXT: ld.param.u32 %r{{[0-9]+}}, {{\[}}[[BASE]]+4{{\]}}
+; PTX-NOT: ld.param.u32 {{%r[0-9]+}}, [{{%rd[0-9]+}}]
+; PTX: ld.param.u32 [[value:%r[0-9]+]], [{{%rd[0-9]+}}+4]
   store i32 %0, i32* %output, align 4
+; PTX-NEXT: st.global.u32 [{{%rd[0-9]+}}], [[value]]
   ret void
 }

EMAIL PREFERENCES
  http://reviews.llvm.org/settings/panel/emailpreferences/
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D10322.27334.patch
Type: text/x-patch
Size: 1621 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20150608/f468204a/attachment.bin>