[LLVMbugs] [Bug 16333] New: NVPTX fails to generate loads/stores with literal offsets

[bugzilla-daemon at llvm.org]

http://llvm.org/bugs/show_bug.cgi?id=16333

            Bug ID: 16333
           Summary: NVPTX fails to generate loads/stores with literal
                    offsets
           Product: tools
           Version: trunk
          Hardware: PC
                OS: Linux
            Status: NEW
          Severity: normal
          Priority: P
         Component: llc
          Assignee: unassignedbugs at nondot.org
          Reporter: fwinter at jlab.org
                CC: llvmbugs at cs.uiuc.edu
    Classification: Unclassified

Created attachment 10680
  --> http://llvm.org/bugs/attachment.cgi?id=10680&action=edit
Repro

NVPTX seems to fail to generate loads/stores with literal offsets. Instead it
issues separate offset calculations. I attach a kernel cmplx_mul_ir_builder.ll
for which libnvvm produces the following code:

        shl.b32         %r12, %r4, 2;
        add.s32         %r13, %r3, %r12;
        add.s32         %r14, %r2, %r12;
        ld.global.f32   %f1, [%r14+16384];
        ld.global.f32   %f2, [%r13];
        mul.f32         %f3, %f1, %f2;
        ld.global.f32   %f4, [%r14];
        ld.global.f32   %f5, [%r13+16384];
        fma.rn.f32      %f6, %f4, %f5, %f3;
        mul.f32         %f7, %f1, %f5;
        mul.f32         %f8, %f4, %f2;
        sub.f32         %f9, %f8, %f7;
        add.s32         %r15, %r1, %r12;
        st.global.f32   [%r15], %f9;
        st.global.f32   [%r15+16384], %f6;
        ret;

whereas NVPTX stumbles to

        add.s32         %r6, %r1, 4096;
        mul.wide.s32    %rl4, %r1, 4;
        add.s64         %rl5, %rl2, %rl4;
        ld.f32  %f1, [%rl5];
        mul.wide.s32    %rl6, %r6, 4;
        add.s64         %rl7, %rl2, %rl6;
        ld.f32  %f2, [%rl7];
        add.s64         %rl8, %rl3, %rl4;
        ld.f32  %f3, [%rl8];
        add.s64         %rl9, %rl3, %rl6;
        ld.f32  %f4, [%rl9];
        mul.f32         %f5, %f4, %f1;
        fma.rn.f32      %f6, %f3, %f2, %f5;
        mul.f32         %f7, %f3, %f1;
        neg.f32         %f8, %f4;
        fma.rn.f32      %f9, %f8, %f2, %f7;
        add.s64         %rl10, %rl1, %rl4;
        st.f32  [%rl10], %f9;
        add.s64         %rl11, %rl1, %rl6;
        st.f32  [%rl11], %f6;
        ret;

I cannot tell if this is due to opt passes included in the closed source
version. However, this code compiles to a higher register usage in the final
code, thus a lower performance.

System: CUDA 5.5 (pre-release), sm_35

-- 
You are receiving this mail because:
You are on the CC list for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20130615/10764a9a/attachment.html>