@woct0rdho Awesome, thanks for explanation!
While we're at this, I've been wondering about this specific thing..

For context:
Initially when working on the previous mentioned RingFormer vocoder arch in my fork, which uses conformer:
( chain: vocoder <-> conformer <-> ring attention <-> triton );
https://github.com/codename0og/codename-rvc-fork-4/blob/main/rvc/lib/algorithm/conformer/conformer.py

I had shared memory issues and yeah.. I figured out it's due to RingAttention.. so what I did is simply enable these:
source
Like so:

@triton.autotune(
    configs=[
        triton.Config(
            {"BLOCK_M": 128, "BLOCK_N": 128, "SEQUENCE_PARALLEL": False},
            num_warps=8,
            num_stages=1,
            pre_hook=init_to_zero("DQ"),
        ),
        triton.Config(
            {"BLOCK_M": 128, "BLOCK_N": 128, "SEQUENCE_PARALLEL": True},
            num_warps=8,
            num_stages=1,
            pre_hook=init_to_zero("DQ"),
        ),
        # Other configs seem to give wrong results when seqlen_q % 128 != 0, disabling them for now
        # # Kernel is buggy (give wrong result) if we set BLOCK_m=128, BLOCK_n=64, num_warps=*4*
        # triton.Config({"BLOCK_M": 128, "BLOCK_N": 64, "SEQUENCE_PARALLEL": False}, num_warps=8, num_stages=1, pre_hook=init_to_zero('DQ')),
        # triton.Config({"BLOCK_M": 128, "BLOCK_N": 64, "SEQUENCE_PARALLEL": True}, num_warps=8, num_stages=1, pre_hook=init_to_zero('DQ')),
        
        # Codename;0 ~ 2 configs below seem to work just fine for RingFormer vocoder. ( Tested on rtx 3060 12gb. )
        # If you spotted any issues related to that, please open up an issue and let me know!
        triton.Config({"BLOCK_M": 64, "BLOCK_N": 64, "SEQUENCE_PARALLEL": False}, num_warps=4, num_stages=1, pre_hook=init_to_zero('DQ')),
        triton.Config({"BLOCK_M": 64, "BLOCK_N": 64, "SEQUENCE_PARALLEL": True}, num_warps=4, num_stages=1, pre_hook=init_to_zero('DQ')),
    ],
    key=["CACHE_KEY_SEQLEN_Q", "CACHE_KEY_SEQLEN_K", "BIAS_TYPE", "IS_CAUSAL", "BLOCK_HEADDIM"],
)

Thanks to this tweak, all goes butter-smooth on RTX 30xx + ( Or so it seems so far?),
However my testers, those with rtx 20xx ( And I believe gtx too ) encounter:

triton.runtime.errors.OutOfResources: out of resource: shared memory, Required: 98304, Hardware limit: 65536.

Not quite sure if there's anything that can be done?
Might have been foolish of me but, I attempted adding this ( No luck, didn't help.):

        triton.Config({"BLOCK_M": 32, "BLOCK_N": 32, "SEQUENCE_PARALLEL": False}, num_warps=2, num_stages=1, pre_hook=init_to_zero('DQ')),
        triton.Config({"BLOCK_M": 32, "BLOCK_N": 32, "SEQUENCE_PARALLEL": True}, num_warps=2, num_stages=1, pre_hook=init_to_zero('DQ')),

So.. Is it simply a 6 or 8 gig gpu users are locked out kind of scenario?
Or is it some specific bug / below ampere-ampere-cards issue?

Edit: or could it be related to usage of fp32? ( vs tf32/bf16 ; as for me and other users )
if so, that'd make sense.. there's a fallback I added for below ampere archs 👀

Thank you in advance 🙏

Possible Pascal GPU support #133

Description

Metadata

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

Issue actions