[opt] Added llvm::SeparateConstOffsetFromGEPPass() for shared_memory optimizations

[jim19930609]

Related issue = #5472

There's an update for SeparateConstOffsetFromGEPPass() from LLVM10 to LLVM15, supported more patterns and situations.

To enhance it in LLVM10, three dependent passes were added manually, namely LoopStrengthReducePass, IndVarSimplifyPass, and EarlyCSEPass. Once SeparateConstOffsetFromGEPPass() doesn't work as expected for your code, try recompiling Taichi with LLVM15 and likely it's gonna work.

Example:

T : ti.i64 = 40
t_group : ti.i64 = 4 
b_group : ti.i64 = 8 
block_dim : ti.i64 = 192 
Tmax : ti.i64 = 768 

x = ti.field(ti.f32, shape=1)

@ti.kernel
def test():
    ti.loop_config(block_dim=block_dim)
    for t_block in ti.ndrange(T // t_group):
        t = t_block * t_group
        w_pad = ti.simt.block.SharedArray((Tmax,), ti.f32)
    
        ti.simt.block.sync()

        for u in range(0, t+1):
            for bi in ti.static(range(b_group)):
                for i in ti.static(range(t_group)):
                    x[0] = w_pad[(u+bi+i)]

        ti.simt.block.sync()

test()

Before Patch:
[LLVM]

[PTX]

After Patch:
[LLVM]

[PTX]

[netlify]

✅ Deploy Preview for docsite-preview ready!

Name	Link
🔨 Latest commit	5689671
🔍 Latest deploy log	https://app.netlify.com/sites/docsite-preview/deploys/62e371e2b6676a000807f728
😎 Deploy Preview	https://deploy-preview-5494--docsite-preview.netlify.app
📱 Preview on mobile	Toggle QR Code... Use your smartphone camera to open QR code link.

---

To edit notification comments on pull requests, go to your Netlify site settings.

[ailzhang]

Thanks!!