Possible Pascal GPU support

[xeghesoopn]

Do you think this can be made to work with the patch to add Pascal GPU support found here (https://github.com/sasha0552/pascal-pkgs-ci)? Having a real hard time building it from source, and wondering whether I'm barking up the wrong tree.

[woct0rdho]

This looks interesting! But I don't have the hardware to test it. If you successfully build it and run the unit tests on a Pascal GPU, please tell me and I'll merge the patches.

I think we can focus on Triton 3.2 on Pascal, because Triton 3.3 dropped support for Turing. You can clone this repo, checkout branch release/3.2.x-windows, apply the patches here: https://github.com/sasha0552/pascal-pkgs-ci/tree/main/patches/triton-lang/triton/release/3.0.x (that's the same as 3.2.x), then build following the instructions here: https://github.com/woct0rdho/triton-windows/blob/readme/BUILD.md

Feel free to ask if you see any error.

[codename0og]

@woct0rdho @xeghesoopn Any news in this particular case?
Asking because RingFormer neural vocoder architecture uses Ring attention which depends on triton,
and as much as I already assessed Turing gpus being the minimum to get it working..
Am wondering if pascal users could see some love 👀

[woct0rdho]

@codename0og If you have a Pascal GPU, you can apply the patches and test it. Feel free to ask if you see any error.

[FurkanGozukara]

@codename0og If you have a Pascal GPU, you can apply the patches and test it. Feel free to ask if you see any error.

@woct0rdho do we have install link that will make flash attention work on rtx 2000 1000 gpus? at least partially?

[woct0rdho]

@FurkanGozukara You can do pip install flash-attn-triton and it should work on Turing, see https://github.com/rationalism/flash-attn-triton

Once Triton works on Pascal, flash-attn-triton should also work

[FurkanGozukara]

@FurkanGozukara You can do pip install flash-attn-triton and it should work on RTX 20xx (Turing), see https://github.com/rationalism/flash-attn-triton

Once Triton works on GTX 10xx (Pascal), flash-attn-triton should also work

awesome ty so much

[codename0og]

@woct0rdho
I actually am running stuff on ampere ( 3060 ) so sadly no, don't have any pascal gpus laying around neither any potential testers, unfortunately

[blap]

Sorry, I just created a PR for FP8 and BF16 (#138) without knowing about sasha0552/pascal-pkgs-ci. Maybe my PR can support other hardware as well.

[blap]

https://docs.nvidia.com/cuda/cuda-toolkit-release-notes/index.html

2.6.2. Deprecated Architectures
Architecture support for Maxwell, Pascal, and Volta is considered feature-complete. Offline compilation and library support for these architectures have been removed in CUDA Toolkit 13.0 major version release. The use of CUDA Toolkits through the 12.x series to build applications for these architectures will continue to be supported, but newer toolkits will be unable to target these architectures.

[woct0rdho]

I'm working on #138 , which is about supporting dtype conversions with fp8 and bf16. It's not directly related to supporting Pascal in general.

In the meanwhile, you can help test if the modification to third_party/nvidia/backend/compiler.py solves some problems on Pascal. It disables some passes that looks not applicable on Pascal.

Even if we don't implement the dtype conversions, I think a short-term goal is to make it possible to run SageAttention 1 on Pascal, see thu-ml/SageAttention#66 . SageAttention uses int8, and it does not directly require fp16/bf16/fp8. Consumer-grade Pascal GPUs such as GTX 10xx have high-performance int8 cores, but they don't have high-performance fp16 cores. You can install SageAttention 1 (the Triton version) and test it using my test_sageattn.py.

[codename0og]

@woct0rdho
This makes me wonder ( pardon my ignorance in advance, I am still very new to all of it )
What is the current status of bf16 in triton? Any consequences or drawbacks?

For context:
My codebase uses a RingFormer vocoder which utilizes RingAttention and that one depends on triton.
Training's done in amp tf32;bf16

Is there any potential issues from going for bf16, at current stage?

[woct0rdho]

@codename0og There should be no problem running bf16 on RTX 30xx and newer. Training and inferencing in bf16 are pretty mature, with less time and memory compared to fp32, and less accuracy loss compared to fp16. But RTX 20xx and older don't support bf16. Feel free to ask if you have more specific questions.

[codename0og]

@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 🙏

[woct0rdho]

@codename0og What if you set BLOCK_M and BLOCK_N to 16? Basically you need smaller block sizes on GPUs with smaller shared memory. But the block size must be a power of 2 and at least 16.