Building on Jetson AGX Xavier Development Kit fails

[g588928812]

Hi,

i am trying to build bitsandbytes on a Nvidia Jetson AGX Xavier Kit, but it fails, not finding emmintrin.h:

/home/g/bitsandbytes# CUDA_VERSION=114 make cuda11x_nomatmul

ENVIRONMENT
============================
CUDA_VERSION: 114
============================
NVCC path: /usr/local/cuda/bin/nvcc
GPP path: /usr/bin/g++ VERSION: g++ (Ubuntu 9.4.0-1ubuntu1~20.04.1) 9.4.0
CUDA_HOME: /usr/local/cuda
CONDA_PREFIX:
PATH: /usr/local/sbin:/usr/local/bin:/usr/sbin:/usr/bin:/sbin:/bin:/usr/games:/usr/local/games:/snap/bin:/usr/local/cuda/bin/
LD_LIBRARY_PATH:
============================
/usr/local/cuda/bin/nvcc -gencode arch=compute_50,code=sm_50 -gencode arch=compute_52,code=sm_52 -gencode arch=compute_60,code=sm_60 -gencode arch=compute_61,code=sm_61 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_72,code=sm_72 -gencode arch=compute_75,code=sm_75 -gencode arch=compute_80,code=sm_80 -gencode arch=compute_86,code=sm_86 -gencode arch=compute_35,code=sm_35 -gencode arch=compute_37,code=sm_37 -Xcompiler '-fPIC' --use_fast_math -Xptxas=-v -dc /home/g/bitsandbytes/csrc/ops.cu /home/g/bitsandbytes/csrc/kernels.cu -I /home/g/sse2neon -I /usr/local/cuda/include -I /home/g/bitsandbytes/csrc -I /include -I /home/g/bitsandbytes/include -L /usr/local/cuda/lib64 -lcudart -lcublas -lcublasLt -lcurand -lcusparse -L /lib --output-directory /home/g/bitsandbytes/build -D NO_CUBLASLT
nvcc warning : The 'compute_35', 'compute_37', 'compute_50', 'sm_35', 'sm_37' and 'sm_50' architectures are deprecated, and may be removed in a future release (Use -Wno-deprecated-gpu-targets to suppress warning).
In file included from /home/g/bitsandbytes/include/BinSearch.h:5,
from /home/g/bitsandbytes/csrc/ops.cu:10:
/home/g/bitsandbytes/include/SIMD.h:32:10: fatal error: emmintrin.h: No such file or directory
32 | #include <emmintrin.h>
| ^~~~~~~~~~~~~
compilation terminated.
make: *** [Makefile:83: cuda11x_nomatmul] Error 1

Did a bit of research and, not knowing what i am doing, I changed SMID.h to include sse2neon.h instead of emmintrin.h. NOW it fails again, catastrophically, not finding builtin functions:

/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(38): error: identifier "__Int8x8_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(39): error: identifier "__Int16x4_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(40): error: identifier "__Int32x2_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(41): error: identifier "__Int64x1_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(42): error: identifier "__Float16x4_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(43): error: identifier "__Float32x2_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(44): error: identifier "__Poly8x8_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(45): error: identifier "__Poly16x4_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(46): error: identifier "__Uint8x8_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(47): error: identifier "__Uint16x4_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(48): error: identifier "__Uint32x2_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(49): error: identifier "__Float64x1_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(50): error: identifier "__Uint64x1_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(51): error: identifier "__Int8x16_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(52): error: identifier "__Int16x8_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(53): error: identifier "__Int32x4_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(54): error: identifier "__Int64x2_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(55): error: identifier "__Float16x8_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(56): error: identifier "__Float32x4_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(57): error: identifier "__Float64x2_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(58): error: identifier "__Poly8x16_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(59): error: identifier "__Poly16x8_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(60): error: identifier "__Poly64x2_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(61): error: identifier "__Poly64x1_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(62): error: identifier "__Uint8x16_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(63): error: identifier "__Uint16x8_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(64): error: identifier "__Uint32x4_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(65): error: identifier "__Uint64x2_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(67): error: identifier "__Poly8_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(68): error: identifier "__Poly16_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(69): error: identifier "__Poly64_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(70): error: identifier "__Poly128_t" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(72): error: identifier "__fp16" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(795): error: identifier "__builtin_aarch64_saddlv8qi" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(802): error: identifier "__builtin_aarch64_saddlv4hi" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(809): error: identifier "__builtin_aarch64_saddlv2si" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(816): error: identifier "__builtin_aarch64_uaddlv8qi" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(824): error: identifier "__builtin_aarch64_uaddlv4hi" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(832): error: identifier "__builtin_aarch64_uaddlv2si" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(840): error: identifier "__builtin_aarch64_saddl2v16qi" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(847): error: identifier "__builtin_aarch64_saddl2v8hi" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(854): error: identifier "__builtin_aarch64_saddl2v4si" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(861): error: identifier "__builtin_aarch64_uaddl2v16qi" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(869): error: identifier "__builtin_aarch64_uaddl2v8hi" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(877): error: identifier "__builtin_aarch64_uaddl2v4si" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(885): error: identifier "__builtin_aarch64_saddwv8qi" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(892): error: identifier "__builtin_aarch64_saddwv4hi" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(899): error: identifier "__builtin_aarch64_saddwv2si" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(906): error: identifier "__builtin_aarch64_uaddwv8qi" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(914): error: identifier "__builtin_aarch64_uaddwv4hi" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(922): error: identifier "__builtin_aarch64_uaddwv2si" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(930): error: identifier "__builtin_aarch64_saddw2v16qi" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(937): error: identifier "__builtin_aarch64_saddw2v8hi" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(944): error: identifier "__builtin_aarch64_saddw2v4si" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(951): error: identifier "__builtin_aarch64_uaddw2v16qi" is undefined
/usr/lib/gcc/aarch64-linux-gnu/9/include/arm_neon.h(959): error: identifier "__builtin_aarch64_uaddw2v8hi" is undefined

SETUP:

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2022 NVIDIA Corporation
Built on Sun_Oct_23_22:16:07_PDT_2022
Cuda compilation tools, release 11.4, V11.4.315
Build cuda_11.4.r11.4/compiler.31964100_0

Flashed using JetPack 5.1 (Ubuntu 20.04)

R35 (release), REVISION: 2.1, GCID: 32413640, BOARD: t186ref, EABI: aarch64, DATE: Tue Jan 24 23:38:33 UTC 2023
Linux ubuntu 5.10.104-tegra #1 SMP PREEMPT Tue Jan 24 15:09:44 PST 2023 aarch64 aarch64 aarch64 GNU/Linux

full_output_nvcc-verbose.txt

Any help would be greatly appreciated, thank you!

[janrinze]

Same issue here. Apparently nvcc does not support NEON.

[g588928812]

OK it works now, I have included the arm_neon intrinsics and moved the parts using those intructions to a separate .cpp file.

The compiled library loads correctly in python, does not complain but produces wrong results (I think).

Here's what happens when quantitizing a very small model (code from https://huggingface.co/blog/hf-bitsandbytes-integration):

import torch
import torch.nn as nn

import bitsandbytes as bnb
from bitsandbytes.nn import Linear8bitLt

fp16_model = nn.Sequential(
    nn.Linear(64, 64),
    nn.Linear(64, 64)
)

torch.save(fp16_model.state_dict(), "model.pt")

int8_model = nn.Sequential(
    Linear8bitLt(64, 64, has_fp16_weights=False),
    Linear8bitLt(64, 64, has_fp16_weights=False)
)

int8_model.load_state_dict(torch.load("model.pt"))

int8_model[0].weight

Output (=before quantitization):

Parameter containing:
Parameter(Int8Params([[-0.0271,  0.0165,  0.0010,  ...,  0.0662,  0.1013, -0.1198],
            [-0.0512,  0.0790,  0.0168,  ...,  0.0641,  0.0217, -0.0149],
            [ 0.0428, -0.0957,  0.0995,  ...,  0.0269,  0.1067,  0.0185],
            ...,
            [ 0.0833, -0.0097,  0.0922,  ..., -0.0274,  0.0309, -0.0728],
            [-0.0933, -0.0087,  0.0862,  ..., -0.1061, -0.0052, -0.1229],
            [ 0.0371,  0.0090,  0.1147,  ...,  0.0673,  0.0801, -0.0073]]))

Now calling int8_model = int8_model.to(0) to 8bitquantitize and again looking at int8_model[0].weight:

Parameter containing:
Parameter(Int8Params([[  0,  17,   1,  ...,  67, 103,   0],
            [  0,  81,  17,  ...,  65,  22,   0],
            [ 44,   0, 101,  ...,  27, 109,  19],
            ...,
            [ 85,   0,  94,  ...,   0,  32,   0],
            [  0,   0,  88,  ...,   0,   0,   0],
            [ 38,   9, 119,  ...,  70,  83,   0]], device='cuda:0',
           dtype=torch.int8))

It seems like all the negative values have been set to zero. Running the same thing in Google Colab works produces negative integers, no zeros.

Any ideas why that is? (Note: I did not modify any of the .cpp/.cu functions, I just moved stuff around)

can you share the modified code? we are facing the same problem, and can debug together.

[janrinze]

OK it works now, I have included the arm_neon intrinsics and moved the parts using those intructions to a separate .cpp file.

The compiled library loads correctly in python, does not complain but produces wrong results (I think).

Here's what happens when quantitizing a very small model (code from https://huggingface.co/blog/hf-bitsandbytes-integration):

import torch
import torch.nn as nn

import bitsandbytes as bnb
from bitsandbytes.nn import Linear8bitLt

fp16_model = nn.Sequential(
    nn.Linear(64, 64),
    nn.Linear(64, 64)
)

torch.save(fp16_model.state_dict(), "model.pt")

int8_model = nn.Sequential(
    Linear8bitLt(64, 64, has_fp16_weights=False),
    Linear8bitLt(64, 64, has_fp16_weights=False)
)

int8_model.load_state_dict(torch.load("model.pt"))

int8_model[0].weight

Output (=before quantitization):

Parameter containing:
Parameter(Int8Params([[-0.0271,  0.0165,  0.0010,  ...,  0.0662,  0.1013, -0.1198],
            [-0.0512,  0.0790,  0.0168,  ...,  0.0641,  0.0217, -0.0149],
            [ 0.0428, -0.0957,  0.0995,  ...,  0.0269,  0.1067,  0.0185],
            ...,
            [ 0.0833, -0.0097,  0.0922,  ..., -0.0274,  0.0309, -0.0728],
            [-0.0933, -0.0087,  0.0862,  ..., -0.1061, -0.0052, -0.1229],
            [ 0.0371,  0.0090,  0.1147,  ...,  0.0673,  0.0801, -0.0073]]))

Now calling int8_model = int8_model.to(0) to 8bitquantitize and again looking at int8_model[0].weight:

Parameter containing:
Parameter(Int8Params([[  0,  17,   1,  ...,  67, 103,   0],
            [  0,  81,  17,  ...,  65,  22,   0],
            [ 44,   0, 101,  ...,  27, 109,  19],
            ...,
            [ 85,   0,  94,  ...,   0,  32,   0],
            [  0,   0,  88,  ...,   0,   0,   0],
            [ 38,   9, 119,  ...,  70,  83,   0]], device='cuda:0',
           dtype=torch.int8))

It seems like all the negative values have been set to zero. Running the same thing in Google Colab works produces negative integers, no zeros.

Any ideas why that is? (Note: I did not modify any of the .cpp/.cu functions, I just moved stuff around)

@g588928812 If you create a fork in your github account you can do the changes there. That way it is possible to collaborate and also to send a pull request.

[janrinze]

is the issue of no negative numbers related to this: pytorch/pytorch#52146 ?

[g588928812]

can you share the modified code? we are facing the same problem, and can debug together.

sure. here's the fork: https://github.com/g588928812/bitsandbytes_jetsonX

[g588928812]

is the issue of no negative numbers related to this: pytorch/pytorch#52146 ?

i'm using pytorch 1.14 and this should be have been fixed with 1.8.1 already if i understand this correctly

[g588928812]

is the issue of no negative numbers related to this: pytorch/pytorch#52146 ?

you were right! I systematically replaced all chars with in8_t and it works now, it was somewhere in kernels.cu. will find out which change exactly did it and update the repository later

thanks for your repo

the test with torch 2.0 has a similar effect.

now that you have located the kernels.cu, i will also try to modify it

[janrinze]

would -fsigned-char suffice?
only on the affected c/cpp file. however, -fsigned-char does make code slower on some parts. So use with caution. Best way is always to use int8_t if that's how the math is supposed to work. Using char in math is probably ambiguous.

[g588928812]

thanks for your repo

the test with torch 2.0 has a similar effect.

now that you have located the kernels.cu, i will also try to modify it

it's done, repository updated. but i dont know yet if the rest of the library works fine, i will check the pytests and see what they produce

[janrinze]

Support for Apple silicon #252 shows another Aarch64 approach. Would be a good idea to merge these efforts.

[rickardp]

Another thing that needs looking into is building proper platform-specific wheels. I've started investigating this in #257, but it is not 100% working yet. I don't know if someone else have started looking into it. As I don't have access to CUDA on Linux currently (my NVIDIA PC is hijacked by my kids for gaming and it runs Win11 right now :) ) I've also tried setting it up on GitHub pipelines.
+1 on janrinze, let's coordinate efforts. I am happy to adapt to whatever is preferred and pitch in where needed! I can probably contribute the most around the build and packaging right now, as I've never used the Neon instructions (though I have coded my fair share of SSE back in the days).

[rickardp]

I think I might have solved this in #257, but I have no hardware to test it on so I can't verify. Wheel is built by https://github.com/rickardp/bitsandbytes/actions/runs/4653237487/jobs/8233937589 (go to Summary, Artifacts then download the bdist_wheel zip and get the correct one).

I had to move some #ifdefs and includes around. Then it seems nvcc doesn't like Neon intrinsics, so I had to compile the Cuda version without Neon support. If anyone has the hardware to try, please check out this build