diff options
Diffstat (limited to 'llvm')
| -rw-r--r-- | llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp | 5 | ||||
| -rw-r--r-- | llvm/test/CodeGen/NVPTX/bug21465.ll | 5 |
2 files changed, 8 insertions, 2 deletions
diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp index 7fdfdbfb83f..a6466687bc7 100644 --- a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp @@ -165,6 +165,11 @@ void NVPTXPassConfig::addIRPasses() { 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. diff --git a/llvm/test/CodeGen/NVPTX/bug21465.ll b/llvm/test/CodeGen/NVPTX/bug21465.ll index c375cf8d580..2eae41f73a0 100644 --- a/llvm/test/CodeGen/NVPTX/bug21465.ll +++ b/llvm/test/CodeGen/NVPTX/bug21465.ll @@ -14,9 +14,10 @@ entry: ; 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 } |
