BitsandBytes Enablement on ROCm #1207
hipify the csrc repo
hipify pythoninterface
copy from agrocylo
hipify cuparse and cublas calls
fix compile error and Makefile
fixed runtime error (low accuracy)
FIX LOW ACCURACY
Update README.md
add benchmarks
Update README.md
First draft, getting error
Small transform fix, still errors on igemm
create HIP_ENVIRONMENT variable
Skip failing tests on rocm
Add default value for HIP_ENVIRONMENT
Merge pull request #1 from ROCmSoftwarePlatform/skip_rocm_failing_tests
skip failing triton tests on rocm
Merge pull request #2 from ROCmSoftwarePlatform/skip_triton
Enable col to row transformation
Add make functions for row to col transformation
Update get_transform_buffer for row to col in HIP
Update igemmlt for col format
Unskip test_igemmlt_int on ROCm
Update igemmlt_int test for col inputs
Skip transpose igemmlt test on ROCm
Revert "Update igemmlt_int test for col inputs"
Return nvidia_transform from transform for HIP
Fix syntax error
Add comment for shape change
Enable nvidia_transform tests
Merge branch 'fix_igemmlt_int' of https://github.com/pnunna93/bitsandβ¦
Enable igemmlt_half tests
Revert col32 check in nvidia_transform test
Merge pull request #3 from pnunna93/fix_igemmlt_int
Merge remote-tracking branch 'upstream/main' into IFU-master-2024-01-24
Update README.md
Update hip files with upstream changes
Skip failing tests for now
Merge pull request #4 from ROCm/IFU-master-2024-01-24
ops.hip: adapt to enum naming changes in ROCm/hipBLASLt@95131d6 and Rβ¦
fix wmma api parity
hipify wmma datatype
Enable estimate quantile tests
Merge pull request #5 from iiisak/rocm_enabled
Merge pull request #7 from ROCm/fix_estimate_quantiles
Merge pull request #6 from ROCm/rocwmma_merge
Enable transpose flag for row to col transform
Update descriptors for transpose flag
revert nvidia_transform to transform
Merge pull request #8 from ROCm/enable_transform_with_transpose
Merge pull request #9 from ROCm/rocm_enabled_fix_bfloat16
remove blocksize 64 on rocm
remove block size 64 and enable remaining tests
Fix cuda build errors
remove workspace in igemmlt
Enabled igemmlt in matmul
Fix shape issue in transform function
Enable igemmlt int8 output
Add col format for extract outliers
Enable dequant_mm
Enable matmullt tests
Enabled linear_serialization tests
fix error with dequant_mm change
Enable extract outliers test
Enable test overflow
Skip overflow and linear serialization for now
Merge pull request #10 from ROCm/remove_blocksize_64
Merge pull request #11 from ROCm/fix_cuda_build_errs
Merge pull request #12 from ROCm/igemm_workspace
Merge pull request #13 from ROCm/enable_matmul
Merge pull request #14 from ROCm/fix_gemv_4bit
Update skip comment
Merge pull request #15 from ROCm/gemv_skip_comment
Merge remote-tracking branch 'upstream/main' into IFU-master-2024-03-28
Merge pull request #19 from ROCm/updated_readme
Update README.md
Merge branch 'rocm_enabled' into IFU-master-2024-03-28
fix PEP errors
Fix typos
Merge branch 'IFU-master-2024-03-28' of https://github.com/ROCm/bitsaβ¦
Fix formatting in README file
Update gpu arch setting
Add ROCM_PATH variable
Add HIP_VERSION variable
Add BNB_HIP_VERSION variable
Update supports igemmlt based on HIP version
Skip failing tests based on HIP version
pre-commit fixes
Update README file
Update default arch list
update readme
Merge pull request #17 from ROCm/IFU-master-2024-03-28
Merge remote-tracking branch 'TD_BnB/multi-backend-refactor' into devβ¦
update igemmlt for hip
Update mm_dequant for hip
Update transform function for hip
Merge pull request #21 from ROCm/rocm_enabled_arch_detect
Merge branch 'rocm_enabled' into device_abstraction
update extract_outliers, quantize_4bit, dequantize_4bit
minor fixes for extract_outliers
update blocksizes for quantize and dequantize
update reg expression for detecting arch
linter updates
Merge branch 'device_abstraction' into cl/update-device-abs
Merge pull request #23 from ROCm/cl/update-device-abs
Merge remote-tracking branch 'upstream/multi-backend-refactor' into dβ¦
skip linear no igemmlt test
Remove archive functional file
Sync README with upstream
Remove bnb_accuracy file
Remove cuda_setup
Remove test_delete_later.c
Sync with upstream
Sync files with upstream
Fix lint errors
Exclude hip files from typo checks
update ops.hip
Merge pull request #27 from ROCm/dev_abs_IFU
Add install steps for ROCm
Fix lint error
Merge pull request #28 from ROCm/dev_abs_add_install_steps
| 419 | 444 | quant_type: Literal["fp4", "nf4"] = "fp4", | |
| 420 | 445 | quant_storage=torch.uint8, | |
| 421 | 446 | ) -> Tuple[torch.Tensor, QuantState]: | |
| 447 | if blocksize is None: | ||
| 448 | blocksize = 64 if not HIP_ENVIRONMENT else 128 |
matthewdouglas1 year agoβ€ 1Is there a short explanation we can add here to explain why this is the default, and likewise below why 64 is not supported?
pnunna931 year agoπ 2Its because of warpsize difference between AMD and NVIDIA GPUs. I have added comments - 410f499
Issue #149: can Intel Arc GPUs be supported in a similar manner?
Add comments for HIP changes
Thanks! This looks promising.
I will try on both AMD and Intel Arc.
Dear @pnunna93,
thanks to you and your team for the amazing work. We're super excited about this and I'm very happy with what I'm seeing at an initial superficial review.
It would be great to have the AMD runner available relatively soon, otherwise it remains quite messy and work intensive to keep track of the correctness of the various backend implementations. Please let me know what I can do to help and I'll make sure to pull the right strings.
Regarding the review, as communicated in Slack, I have to first focus on wrapping up my deep dive in evaluating tensor-driven dispatch by integration with the PyTorch dispatcher via the torch.library APIs. I don't see any reason to not merge your PR, but I need to take another thorough look and I think it would be helpful for everyone to have clarity on the backend abstraction / dispatch mechanism asap and am therefore prioritizing that; so everyone can then refactor their code to account for that.
In that context, one important question came up:
Our paged optimizers use CUDA unified memory, as described in detail here.
Is that feature available on ROCm devices in one way or another? This would be quite important to understand for my analysis, as the handling of unified memory in relation to PyTorch is one of my last open questions. It's quite a special case, because it's a cornerstone of preventing OOMs in low resource environments -- a key feature for our user group -- and is not implemented/ accounted for in PyTorch and, therefore, we use that feature directly through CUDA related APIs: The underlying CUDA function is cudaMemPrefetchAsync AFAICT.
Thanks π€ and congrats on the great work in this PR, we're super excited about this β€οΈ
Dear @pnunna93 et al,
Unfortunately we're (mostly me alone) quite resource constrained and humbled by the workload associated with the multi-backend-refactor. I just talked with my colleague @younesbelkada about the topic how to best handle the next steps.
We both took a look at this PR and the one from Intel and think that at first glance everything looks really good. At this time, both me and Younes are not in a position to give detailed feedback and I need to focus on concretizing the path forward on how to integrate with the PyTorch dispatcher (tensor driven dispatch, as requested) through the torch.library Python-level APIs. After extensive research and yesterday's consultation with 3 PyTorch devs at Meta that are experts on the topic I need to focus on making this new input concrete.
However, for the purpose of iterative progress (as agreed in our prior conversations), we've decided to already go ahead and merge both the open Intel and AMD branches into multi-backend-refactor, where interested parties can then compile from source and give the new functionality (we're so excited and grateful about this!) a thorough testing.
Once we've made some progress on the torch.library based refactor, I'll next focus on enabling the nightly releases for that branch as well. We're also looking forward to your feedback on the this torch.library / tensor-driven dispatch topic once the code is there on the basis of which to discuss (and refactor the backend specific code towards that new target, after we agreed with all of you that this is the right path).
Among other things, there's also been extensive ongoing work in the background on things like moving BNB to a new independent/non-profit Github org, but under the umbrella of Hugging Face and the support of their infra team for managing the complexities of the CI/CD backend and runners. Also, we're working to make Github runners for the different hardware platforms a reality (thanks for your help on that!).
Thanks again for the good work and active collaboration! β€οΈ π
P.S. Also see this: README: asking for help from volunteer alpha testers
Let us know if you have further thoughts on this and how you think it's best to communicate about this.
Login to write a write a comment.
Overview
This PR introduces bitsandbytes enablement on ROCm for AMD GPUs. It adds hipified versions of CUDA kernels and ops which allow the flow to route bitsandbytes API function calls to use optimized version of HIP kernels for AMD GPUs.
In the multi-backend-refactor branch, there is a proposal to separate various backends to support multiple GPUs/accelerators. The core of bitsandbytes is built on top of PyTorch and decides the API function call of individual GPUs/accelerators based on the device_type of the tensor as highlighted here. ROCm recognizes cuda device type in PyTorch and runs seamlessly without the need to change anything in the application code. Hence, this PR updates cuda backend in bitsandbytes to enable its functionality on ROCm for AMD GPUs. This PR also adds support for ROCm in the cmake build and enables key functionality of bitsandbytes on AMD GPUs.
Summary of Changes
Updated CUDA backend to work seamlessly on ROCm
Integrated HIP environment into bitsandbytes through hipified versions of CUDA kernels and ops
Cmake build updates for ROCm
Enabled key features in bitsandbytes functional and autograd api
Impact
It enables to build and support bitsandbytes on ROCm for AMD GPUs . Bitsandbytes users can port applications smoothly onto AMD gpus as it requires minimal changes on their front. In addition to this, it also ensures that ROCm changes do not affect CUDA environment, thereby not affecting existing CUDA users.
CC: @Titus-von-Koeller @matthewdouglas @arlo-phoenix